CUDA从入门到精通:性能剖析和Visual Profiler

 
 
 
 
 
 

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

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

 

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


 
 
  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. cudaEventRecord(stop,stream); //记录结束事件
  7. cudaEventSynchronize(stop); //事件同步,等待结束事件之前的设备操作均已完成
  8. float elapsedTime;
  9. cudaEventElapsedTime(&elapsedTime,start,stop); //计算两个事件之间时长(单位为ms)


 

 

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

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


 
 
  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. // cudaThreadExit must be called before exiting in order for profiling and
  42. // tracing tools such as Nsight and Visual Profiler to show complete traces.
  43. cudaStatus = cudaThreadExit();
  44. if (cudaStatus != cudaSuccess)
  45. {
  46. fprintf( stderr, "cudaThreadExit failed!");
  47. return 1;
  48. }
  49. for( int i = 0;i<arraySize;i++)
  50. {
  51. if(c[i] != (a[i]+b[i]))
  52. {
  53. printf( "Error in %d\n",i);
  54. }
  55. }
  56. return 0;
  57. }
  58. // Helper function for using CUDA to add vectors in parallel.
  59. cudaError_t addWithCuda(int *c, const int *a, const int *b, size_t size)
  60. {
  61. int *dev_a = 0;
  62. int *dev_b = 0;
  63. int *dev_c = 0;
  64. cudaError_t cudaStatus;
  65. // Choose which GPU to run on, change this on a multi-GPU system.
  66. cudaStatus = cudaSetDevice( 0);
  67. if (cudaStatus != cudaSuccess)
  68. {
  69. fprintf( stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
  70. goto Error;
  71. }
  72. // Allocate GPU buffers for three vectors (two input, one output) .
  73. cudaStatus = cudaMalloc(( void**)&dev_c, size * sizeof( int));
  74. if (cudaStatus != cudaSuccess)
  75. {
  76. fprintf( stderr, "cudaMalloc failed!");
  77. goto Error;
  78. }
  79. cudaStatus = cudaMalloc(( void**)&dev_a, size * sizeof( int));
  80. if (cudaStatus != cudaSuccess)
  81. {
  82. fprintf( stderr, "cudaMalloc failed!");
  83. goto Error;
  84. }
  85. cudaStatus = cudaMalloc(( void**)&dev_b, size * sizeof( int));
  86. if (cudaStatus != cudaSuccess)
  87. {
  88. fprintf( stderr, "cudaMalloc failed!");
  89. goto Error;
  90. }
  91. // Copy input vectors from host memory to GPU buffers.
  92. cudaStatus = cudaMemcpy(dev_a, a, size * sizeof( int), cudaMemcpyHostToDevice);
  93. if (cudaStatus != cudaSuccess)
  94. {
  95. fprintf( stderr, "cudaMemcpy failed!");
  96. goto Error;
  97. }
  98. cudaStatus = cudaMemcpy(dev_b, b, size * sizeof( int), cudaMemcpyHostToDevice);
  99. if (cudaStatus != cudaSuccess)
  100. {
  101. fprintf( stderr, "cudaMemcpy failed!");
  102. goto Error;
  103. }
  104. cudaEvent_t start,stop;
  105. cudaEventCreate(&start);
  106. cudaEventCreate(&stop);
  107. cudaEventRecord(start, 0);
  108. for( int i = 0;i< 1000;i++)
  109. {
  110. // addKernel_blk<<<size,1>>>(dev_c, dev_a, dev_b);
  111. addKernel_thd<<< 1,size>>>(dev_c, dev_a, dev_b);
  112. }
  113. cudaEventRecord(stop, 0);
  114. cudaEventSynchronize(stop);
  115. float tm;
  116. cudaEventElapsedTime(&tm,start,stop);
  117. printf( "GPU Elapsed time:%.6f ms.\n",tm);
  118. // cudaThreadSynchronize waits for the kernel to finish, and returns
  119. // any errors encountered during the launch.
  120. cudaStatus = cudaThreadSynchronize();
  121. if (cudaStatus != cudaSuccess)
  122. {
  123. fprintf( stderr, "cudaThreadSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
  124. goto Error;
  125. }
  126. // Copy output vector from GPU buffer to host memory.
  127. cudaStatus = cudaMemcpy(c, dev_c, size * sizeof( int), cudaMemcpyDeviceToHost);
  128. if (cudaStatus != cudaSuccess)
  129. {
  130. fprintf( stderr, "cudaMemcpy failed!");
  131. goto Error;
  132. }
  133. Error:
  134. cudaFree(dev_c);
  135. cudaFree(dev_a);
  136. cudaFree(dev_b);
  137. return cudaStatus;
  138. }


 

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

线程并行:

块并行:

 

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

 

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

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

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

 

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

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

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

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

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

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

 

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

 

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

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

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

 

后面Settings就是我们建立会话时的参数配置,不再详述。

 

通过本节,我们应该能对CUDA性能提升有了一些想法,好,下一节我们将讨论如何优化CUDA程序。

  • 0
    点赞
  • 2
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值