您的位置: 首页  >  文章  >  CUDA之Visual Profiler使用

CUDA之Visual Profiler使用

分类: 文章 • 2024-08-07 15:29:10

入门后的进一步学习的内容，就是如何优化自己的代码。我们前面的例子没有考虑任何性能方面优化，是为了更好地学习基本知识点，而不是其他细节问题。从本节开始，我们要从性能出发考虑问题，不断优化代码，使执行速度提高是并行处理的唯一目的。

测试代码运行速度有很多方法，C语言里提供了类似于SystemTime()这样的API获得系统时间，然后计算两个事件之间的时长从而完成计时功能。在CUDA中，我们有专门测量设备运行时间的API，下面一一介绍。

 

翻开编程手册《CUDA_Toolkit_Reference_Manual》，随时准备查询不懂得API。我们在运行核函数前后，做如下操作：

[cpp] view plain copy
  1. cudaEvent_t start,stop;//事件对象  
  2. cudaEventCreate(&start);//创建事件  
  3. cudaEventCreate(&stop);//创建事件  
  4. cudaEventRecord(start,stream);//记录开始  
  5. myKernel<<<dimg,dimb,size_smem,stream>>>(parameter list);//执行核函数  
  6.   
  7. cudaEventRecord(stop,stream);//记录结束事件  
  8. cudaEventSynchronize(stop);//事件同步，等待结束事件之前的设备操作均已完成  
  9. float elapsedTime;  
  10. cudaEventElapsedTime(&elapsedTime,start,stop);//计算两个事件之间时长（单位为ms）  


 

 

核函数执行时间将被保存在变量elapsedTime中。通过这个值我们可以评估算法的性能。下面给一个例子，来看怎么使用计时功能。

前面的例子规模很小，只有5个元素，处理量太小不足以计时，下面将规模扩大为1024，此外将反复运行1000次计算总时间，这样估计不容易受随机扰动影响。我们通过这个例子对比线程并行和块并行的性能如何。代码如下：

[cpp] view plain copy
  1. #include "cuda_runtime.h"  
  2. #include "device_launch_parameters.h"  
  3. #include <stdio.h>  
  4. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size);  
  5. __global__ void addKernel_blk(int *c, const int *a, const int *b)  
  6. {  
  7.     int i = blockIdx.x;  
  8.     c[i] = a[i]+ b[i];  
  9. }  
  10. __global__ void addKernel_thd(int *c, const int *a, const int *b)  
  11. {  
  12.     int i = threadIdx.x;  
  13.     c[i] = a[i]+ b[i];  
  14. }  
  15. int main()  
  16. {  
  17.     const int arraySize = 1024;  
  18.     int a[arraySize] = {0};  
  19.     int b[arraySize] = {0};  
  20.     for(int i = 0;i<arraySize;i++)  
  21.     {  
  22.         a[i] = i;  
  23.         b[i] = arraySize-i;  
  24.     }  
  25.     int c[arraySize] = {0};  
  26.     // Add vectors in parallel.  
  27.     cudaError_t cudaStatus;  
  28.     int num = 0;  
  29.     cudaDeviceProp prop;  
  30.     cudaStatus = cudaGetDeviceCount(&num);  
  31.     for(int i = 0;i<num;i++)  
  32.     {  
  33.         cudaGetDeviceProperties(&prop,i);  
  34.     }  
  35.     cudaStatus = addWithCuda(c, a, b, arraySize);  
  36.     if (cudaStatus != cudaSuccess)   
  37.     {  
  38.         fprintf(stderr, "addWithCuda failed!");  
  39.         return 1;  
  40.     }  
  41.   
  42.     // cudaThreadExit must be called before exiting in order for profiling and  
  43.     // tracing tools such as Nsight and Visual Profiler to show complete traces.  
  44.     cudaStatus = cudaThreadExit();  
  45.     if (cudaStatus != cudaSuccess)   
  46.     {  
  47.         fprintf(stderr, "cudaThreadExit failed!");  
  48.         return 1;  
  49.     }  
  50.     for(int i = 0;i<arraySize;i++)  
  51.     {  
  52.         if(c[i] != (a[i]+b[i]))  
  53.         {  
  54.             printf("Error in %d\n",i);  
  55.         }  
  56.     }  
  57.     return 0;  
  58. }  
  59. // Helper function for using CUDA to add vectors in parallel.  
  60. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)  
  61. {  
  62.     int *dev_a = 0;  
  63.     int *dev_b = 0;  
  64.     int *dev_c = 0;  
  65.     cudaError_t cudaStatus;  
  66.   
  67.     // Choose which GPU to run on, change this on a multi-GPU system.  
  68.     cudaStatus = cudaSetDevice(0);  
  69.     if (cudaStatus != cudaSuccess)   
  70.     {  
  71.         fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");  
  72.         goto Error;  
  73.     }  
  74.     // Allocate GPU buffers for three vectors (two input, one output)    .  
  75.     cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));  
  76.     if (cudaStatus != cudaSuccess)   
  77.     {  
  78.         fprintf(stderr, "cudaMalloc failed!");  
  79.         goto Error;  
  80.     }  
  81.     cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));  
  82.     if (cudaStatus != cudaSuccess)   
  83.     {  
  84.         fprintf(stderr, "cudaMalloc failed!");  
  85.         goto Error;  
  86.     }  
  87.     cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));  
  88.     if (cudaStatus != cudaSuccess)   
  89.     {  
  90.         fprintf(stderr, "cudaMalloc failed!");  
  91.         goto Error;  
  92.     }  
  93.     // Copy input vectors from host memory to GPU buffers.  
  94.     cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);  
  95.     if (cudaStatus != cudaSuccess)   
  96.     {  
  97.         fprintf(stderr, "cudaMemcpy failed!");  
  98.         goto Error;  
  99.     }  
  100.     cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);  
  101.     if (cudaStatus != cudaSuccess)   
  102.     {  
  103.         fprintf(stderr, "cudaMemcpy failed!");  
  104.         goto Error;  
  105.     }  
  106.     cudaEvent_t start,stop;  
  107.     cudaEventCreate(&start);  
  108.     cudaEventCreate(&stop);  
  109.     cudaEventRecord(start,0);  
  110.     for(int i = 0;i<1000;i++)  
  111.     {  
  112. //      addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);  
  113.         addKernel_thd<<<1,size>>>(dev_c, dev_a, dev_b);  
  114.     }  
  115.     cudaEventRecord(stop,0);  
  116.     cudaEventSynchronize(stop);  
  117.     float tm;  
  118.     cudaEventElapsedTime(&tm,start,stop);  
  119.     printf("GPU Elapsed time:%.6f ms.\n",tm);  
  120.     // cudaThreadSynchronize waits for the kernel to finish, and returns  
  121.     // any errors encountered during the launch.  
  122.     cudaStatus = cudaThreadSynchronize();  
  123.     if (cudaStatus != cudaSuccess)   
  124.     {  
  125.         fprintf(stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);  
  126.         goto Error;  
  127.     }  
  128.     // Copy output vector from GPU buffer to host memory.  
  129.     cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);  
  130.     if (cudaStatus != cudaSuccess)   
  131.     {  
  132.         fprintf(stderr, "cudaMemcpy failed!");  
  133.         goto Error;  
  134.     }  
  135. Error:  
  136.     cudaFree(dev_c);  
  137.     cudaFree(dev_a);  
  138.     cudaFree(dev_b);      
  139.     return cudaStatus;  
  140. }  


 

addKernel_blk是采用块并行实现的向量相加操作，而addKernel_thd是采用线程并行实现的向量相加操作。分别运行，得到的结果如下图所示：

线程并行：

CUDA之Visual Profiler使用

块并行：

CUDA之Visual Profiler使用

 

可见性能竟然相差近16倍！因此选择并行处理方法时，如果问题规模不是很大，那么采用线程并行是比较合适的，而大问题分多个线程块处理时，每个块内线程数不要太少，像本文中的只有1个线程，这是对硬件资源的极大浪费。一个理想的方案是，分N个线程块，每个线程块包含512个线程，将问题分解处理，效率往往比单一的线程并行处理或单一块并行处理高很多。这也是CUDA编程的精髓。

 

上面这种分析程序性能的方式比较粗糙，只知道大概运行时间长度，对于设备程序各部分代码执行时间没有一个深入的认识，这样我们就有个问题，如果对代码进行优化，那么优化哪一部分呢？是将线程数调节呢，还是改用共享内存？这个问题最好的解决方案就是利用Visual Profiler。下面内容摘自《CUDA_Profiler_Users_Guide》

“Visual Profiler是一个图形化的剖析工具，可以显示你的应用程序中CPU和GPU的活动情况，利用分析引擎帮助你寻找优化的机会。”

其实除了可视化的界面，NVIDIA提供了命令行方式的剖析命令：nvprof。对于初学者，使用图形化的方式比较容易上手，所以本节使用Visual Profiler。

 

打开Visual Profiler，可以从CUDA Toolkit安装菜单处找到。主界面如下：

CUDA之Visual Profiler使用

我们点击File->New Session，弹出新建会话对话框，如下图所示：

CUDA之Visual Profiler使用

其中File一栏填入我们需要进行剖析的应用程序exe文件，后面可以都不填（如果需要命令行参数，可以在第三行填入），直接Next，见下图：

CUDA之Visual Profiler使用

第一行为应用程序执行超时时间设定，可不填；后面三个单选框都勾上，这样我们分别使能了剖析，使能了并发核函数剖析，然后运行分析器。

点Finish，开始运行我们的应用程序并进行剖析、分析性能。

CUDA之Visual Profiler使用

上图中，CPU和GPU部分显示了硬件和执行内容信息，点某一项则将时间条对应的部分高亮，便于观察，同时右边详细信息会显示运行时间信息。从时间条上看出，cudaMalloc占用了很大一部分时间。下面分析器给出了一些性能提升的关键点，包括：低计算利用率（计算时间只占总时间的1.8%，也难怪，加法计算复杂度本来就很低呀！）；低内存拷贝/计算交叠率（一点都没有交叠，完全是拷贝——计算——拷贝）；低存储拷贝尺寸（输入数据量太小了，相当于你淘宝买了个日记本，运费比实物价格还高！）；低存储拷贝吞吐率（只有1.55GB/s）。这些对我们进一步优化程序是非常有帮助的。

 

我们点一下Details，就在Analysis窗口旁边。得到结果如下所示：

 

CUDA之Visual Profiler使用

通过这个窗口可以看到每个核函数执行时间，以及线程格、线程块尺寸，占用寄存器个数，静态共享内存、动态共享内存大小等参数，以及内存拷贝函数的执行情况。这个提供了比前面cudaEvent函数测时间更精确的方式，直接看到每一步的执行时间，精确到ns。

在Details后面还有一个Console，点一下看看。

CUDA之Visual Profiler使用

这个其实就是命令行窗口，显示运行输出。看到加入了Profiler信息后，总执行时间变长了（原来线程并行版本的程序运行时间只需4ms左右）。这也是“测不准定理”决定的，如果我们希望测量更细微的时间，那么总时间肯定是不准的；如果我们希望测量总时间，那么细微的时间就被忽略掉了。

http://blog.****.net/kkk584520/article/details/9490233

相关推荐