用户
 找回密码
 立即注册
dyanwithu 该用户已被删除
发表于 2013-5-2 22:34:54
5433527
之前在论坛和stackoverflow上看过一些测试GPU浮点运算能力的例子,
例如类似下面的代码,kernel中是大量的MAD运算,没有memory的读取。(这样的代码可以找到很多版本)
有人使用M2090计算,时间为2.16ms,其浮点性能达到1000*1024*30*80/2.16ms = 1.15TFlops vs. 1.33 theoretical peak


问题:我使用的是NVS4200M显卡(2.1),在win7 x64 vs2010 + CUDA5.0环境下的运行时间55ms,浮点性能只有45GFLOPS,而这块显卡的理论峰值大约是100GFLOPS,类似的测试在C2070上运行过,都表现很差。


想请教一下是什么原因导致的实测性能与理论性能相差太大?欢迎讨论,谢谢回复


  1. __global__ void test(float loop, float *out)
  2. {
  3.     register float a=1.0f;
  4.     register float b=1.0f;
  5.     register float c=1.0f;
  6.     register float d=1.0f;
  7.     register float e=1.0f;
  8.     register float f=1.0f;
  9.     register float g=1.0f;
  10.     register float h=1.0f;

  11.     for (float x=0;x<loop;x++)
  12.     {
  13.         a+=x*loop;
  14.         b+=x*loop;
  15.         c+=x*loop;
  16.         d+=x*loop;
  17.         e+=x*loop;
  18.         f+=x*loop;
  19.         g+=x*loop;
  20.         h+=x*loop;

  21.         a+=x*loop;
  22.         b+=x*loop;
  23.         c+=x*loop;
  24.         d+=x*loop;
  25.         e+=x*loop;
  26.         f+=x*loop;
  27.         g+=x*loop;
  28.         h+=x*loop;

  29.         a+=x*loop;
  30.         b+=x*loop;
  31.         c+=x*loop;
  32.         d+=x*loop;
  33.         e+=x*loop;
  34.         f+=x*loop;
  35.         g+=x*loop;
  36.         h+=x*loop;

  37.         a+=x*loop;
  38.         b+=x*loop;
  39.         c+=x*loop;
  40.         d+=x*loop;
  41.         e+=x*loop;
  42.         f+=x*loop;
  43.         g+=x*loop;
  44.         h+=x*loop;

  45.         a+=x*loop;
  46.         b+=x*loop;
  47.         c+=x*loop;
  48.         d+=x*loop;
  49.         e+=x*loop;
  50.         f+=x*loop;
  51.         g+=x*loop;
  52.         h+=x*loop;
  53.     }
  54.     if (out!=NULL) *out=a+b+c+d+e+f+g+h;
  55. }

  56. int main(int argc, char *argv[])
  57. {
  58.     float timestamp;
  59.     cudaEvent_t event_start,event_stop;
  60.     // Initialise
  61.     cudaDeviceReset();
  62.     cudaSetDevice(0);
  63.     cudaThreadSetCacheConfig(cudaFuncCachePreferShared);
  64.     // Allocate and generate buffers
  65.     cudaEventCreate(&event_start);
  66.     cudaEventCreate(&event_stop);
  67.     cudaEventRecord(event_start, 0);
  68.     dim3 threadsPerBlock;
  69.     dim3 blocks;
  70.     threadsPerBlock.x=32;
  71.     threadsPerBlock.y=32;
  72.     threadsPerBlock.z=1;
  73.     blocks.x=1;
  74.     blocks.y=1000;
  75.     blocks.z=1;

  76.     test<<<blocks,threadsPerBlock,0>>>(30,NULL);

  77.     cudaEventRecord(event_stop, 0);
  78.     cudaEventSynchronize(event_stop);
  79.     cudaEventElapsedTime(×tamp, event_start, event_stop);
  80.     printf("Calculated in %f\n", timestamp);
  81. }
复制代码


使用道具 举报 回复
发表于 2013-5-2 22:48:39
如下3个方面可能会影响性能,注意是可能。

(1)建议使用(16,16)的block大小启动贵kernel.
(2)建议使用-O2编译。
(3)建议在启动kernel前单独启动一次,作为热身。而以第二次的计时为准。

请尝试并重新给出结果。
使用道具 举报 回复 支持 反对
发表于 2013-5-2 23:00:24
横扫千军 发表于 2013-5-2 22:48
如下3个方面可能会影响性能,注意是可能。

(1)建议使用(16,16)的block大小启动贵kernel.

谢谢您的回复,按照您的建议进行了以下的测试。

1. (16,16)的block运行时间为13.65ms,达到的浮点性能为45GFLOPS
2. 在CUDA C/C++设置Additional Compiler Options -O2,运行时间基本不变
3. 启动两次kernel,以第二次计时,运行时间基本不变

还是想不明白为什么会差这么多?
使用道具 举报 回复 支持 反对
发表于 2013-5-2 23:02:28
dyanwithu 发表于 2013-5-2 23:00
谢谢您的回复,按照您的建议进行了以下的测试。

1. (16,16)的block运行时间为13.65ms,达到的浮点性能为 ...

我也不明白了。坐等他人回复吧。
使用道具 举报 回复 支持 反对
发表于 2013-5-2 23:07:04
横扫千军 发表于 2013-5-2 23:02
我也不明白了。坐等他人回复吧。

谢谢

类似还有这样的代码,我在nvs4200m上测试过多次,改过许多参数,基本运行速度都在10GFLOPS以下,而C2070上也只有50多GFLOPS,与理论峰值相差太远。

而这些kernel本身又是极其简单的,不存在mem-bound,math上也应该能较好发挥GPU性能,但是实测却差的比较多。(不过提供这些代码的网页上倒有许多不错的测试结果,基本都能达到理论峰值的90%)
  1. #define LOOP (100000)
  2. #define BLOCKS (200)
  3. #define THPB (256)

  4. __global__ void new_ker(float *x)
  5. {
  6.   int index = threadIdx.x+blockIdx.x*blockDim.x;
  7.   float a = 1.0f, b = -1.0f;

  8.   for(int i = 0; i < LOOP; i++){
  9.        a = a*b + b;
  10.   }  

  11.   x[index] = a;
  12. }

  13. int main(int argc,char **argv)
  14. {

  15.    //Initializations
  16.    float *x;
  17.    float *dx;
  18.    cudaEvent_t new_start,new_stop;
  19.    float elapsed;
  20.    double gflops;
  21.    x = (float *)malloc(sizeof(float)*THPB*BLOCKS);
  22.    cudaProfilerStart();
  23.    cudaMalloc((void **)&dx,sizeof(float)*THPB);

  24.    //ILP=1  
  25.    cudaEventCreate(&new_start);
  26.    cudaEventCreate(&new_stop);

  27.    printf("Kernel1:\n");
  28.    cudaEventRecord(new_start, 0);

  29.    new_ker<<<BLOCKS,THPB>>>(dx);

  30.    cudaEventRecord(new_stop,0);
  31.    cudaEventSynchronize(new_stop);
  32.    cudaEventElapsedTime(&elapsed,new_start,new_stop);
  33.    
  34.    cudaMemcpy(x,dx,sizeof(float)*THPB*BLOCKS,cudaMemcpyDeviceToHost);

  35.    cudaEventDestroy(new_start);
  36.    cudaEventDestroy(new_stop);
  37.    cudaDeviceReset();
  38.    cudaProfilerStop();

  39.    printf("\t%f\n",elapsed);
  40.    gflops = 2.0e-6 * ((double)(LOOP)*(double)(THPB*BLOCKS)/(double)elapsed);
  41.    printf("\t%f\n",gflops);

  42.    return 0;
  43. }
复制代码
使用道具 举报 回复 支持 反对
发表于 2013-5-2 23:32:06
dyanwithu 发表于 2013-5-2 23:07
谢谢

类似还有这样的代码,我在nvs4200m上测试过多次,改过许多参数,基本运行速度都在10GFLOPS以下,而 ...

LZ您好,您不妨用NVVP跑一下,看看有什么提示没有。再根据NVVP的结果,调整您的代码,或者能有所启示。

祝您好运~
使用道具 举报 回复 支持 反对
发表于 2013-5-3 00:46:59
我用GTX 650 下指令nvcc *.cu -arch=sm_30
float跑出來是
18.876352/542.477690
理論值是812.5GFLOPs
double跑出來是
149.229248/34.309628
使用道具 举报 回复 支持 反对
发表于 2013-5-3 03:57:51
本帖最后由 横扫千军 于 2013-5-3 04:45 编辑
dyanwithu 发表于 2013-5-2 23:07
谢谢

类似还有这样的代码,我在nvs4200m上测试过多次,改过许多参数,基本运行速度都在10GFLOPS以下,而 ...

无法理解了。好奇怪的。

以及,建议第二个kernel进行#pragma unroll下。循环体太小。

建议其他人给出解答。
使用道具 举报 回复 支持 反对
发表于 2013-5-3 10:10:56
横扫千军 发表于 2013-5-3 03:57
无法理解了。好奇怪的。

以及,建议第二个kernel进行#pragma unroll下。循环体太小。

谢谢
第二个kernel原本是带unroll的,但是对结果没有很大的影响。
使用道具 举报 回复 支持 反对
发表于 2013-5-3 10:11:39
iHakka 发表于 2013-5-3 00:46
我用GTX 650 下指令nvcc *.cu -arch=sm_30
float跑出來是
18.876352/542.477690

你的这个结果还不错,改改参数应该能到理论峰值的90%
使用道具 举报 回复 支持 反对
123下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册