用户
 找回密码
 立即注册
lzhbch 该用户已被删除
发表于 2013-9-10 15:11:17
61495
本帖最后由 lzhbch 于 2013-9-10 15:14 编辑

先描述一下自己的环境,win7 x64系统,VS2005, 650M
问题1:
       同样一段代码,我用默认的CUDA 5.0 Runtime工程来跑与我把CUDA sample中的工程替换成同样的代码来跑,执行效率差距很大,不知道是什么原因造成的?sample使用的工程是simpleStreams工程。

首先是在CUDA 5.0 Runtime工程中的运行结果


然后是sample工程中的运行结果


所使用的程序来自CUDA的官方文章:How to Overlap Data Transfers in CUDA C/C++
代码为:
  1. inline
  2. cudaError_t checkCuda(cudaError_t result)
  3. {
  4. #if defined(DEBUG) || defined(_DEBUG)
  5.   if (result != cudaSuccess) {
  6.     fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
  7.     assert(result == cudaSuccess);
  8.   }
  9. #endif
  10.   return result;
  11. }

  12. __global__ void kernel(float *a, int offset)
  13. {
  14.   int i = offset + threadIdx.x + blockIdx.x*blockDim.x;
  15.   float x = (float)i;
  16.   float s = sinf(x);
  17.   float c = cosf(x);
  18.   a[i] = a[i] + sqrtf(s*s+c*c);
  19. }

  20. float maxError(float *a, int n)
  21. {
  22.   float maxE = 0;
  23.   for (int i = 0; i < n; i++) {
  24.     float error = fabs(a[i]-1.0f);
  25.     if (error > maxE) maxE = error;
  26.   }
  27.   return maxE;
  28. }

  29. int main(int argc, char **argv)
  30. {
  31.   const int blockSize = 256, nStreams = 4;
  32.   const int n = 4 * 1024 * blockSize * nStreams;
  33.   const int streamSize = n / nStreams;
  34.   const int streamBytes = streamSize * sizeof(float);
  35.   const int bytes = n * sizeof(float);
  36.    
  37.   int devId = 0;
  38.   if (argc > 1) devId = atoi(argv[1]);

  39.   cudaDeviceProp prop;
  40.   checkCuda( cudaGetDeviceProperties(&prop, devId));
  41.   printf("Device : %s\n", prop.name);
  42.   checkCuda( cudaSetDevice(devId) );
  43.   
  44.   // allocate pinned host memory and device memory
  45.   float *a, *d_a;
  46.   checkCuda( cudaMallocHost((void**)&a, bytes) );      // host pinned
  47.   checkCuda( cudaMalloc((void**)&d_a, bytes) ); // device

  48.   float ms; // elapsed time in milliseconds
  49.   
  50.   // create events and streams
  51.   cudaEvent_t startEvent, stopEvent, dummyEvent;
  52.   cudaStream_t stream[nStreams];
  53.   checkCuda( cudaEventCreate(&startEvent) );
  54.   checkCuda( cudaEventCreate(&stopEvent) );
  55.   checkCuda( cudaEventCreate(&dummyEvent) );
  56.   for (int i = 0; i < nStreams; ++i)
  57.     checkCuda( cudaStreamCreate(&stream[i]) );
  58.   
  59.   // baseline case - sequential transfer and execute
  60.   memset(a, 0, bytes);
  61.   checkCuda( cudaEventRecord(startEvent,0) );
  62.   checkCuda( cudaMemcpy(d_a, a, bytes, cudaMemcpyHostToDevice) );
  63.   kernel<<<n/blockSize, blockSize>>>(d_a, 0);
  64.   checkCuda( cudaMemcpy(a, d_a, bytes, cudaMemcpyDeviceToHost) );
  65.   checkCuda( cudaEventRecord(stopEvent, 0) );
  66.   checkCuda( cudaEventSynchronize(stopEvent) );
  67.   checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
  68.   printf("Time for sequential transfer and execute (ms): %f\n", ms);
  69.   printf("  max error: %e\n", maxError(a, n));

  70.   // asynchronous version 1: loop over {copy, kernel, copy}
  71.   memset(a, 0, bytes);
  72.   checkCuda( cudaEventRecord(startEvent,0) );
  73.   for (int i = 0; i < nStreams; ++i) {
  74.     int offset = i * streamSize;
  75.     checkCuda( cudaMemcpyAsync(&d_a[offset], &a[offset],
  76.                                streamBytes, cudaMemcpyHostToDevice,
  77.                                stream[i]) );
  78.     kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
  79.     checkCuda( cudaMemcpyAsync(&a[offset], &d_a[offset],
  80.                                streamBytes, cudaMemcpyDeviceToHost,
  81.                                stream[i]) );
  82.   }
  83.   checkCuda( cudaEventRecord(stopEvent, 0) );
  84.   checkCuda( cudaEventSynchronize(stopEvent) );
  85.   checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
  86.   printf("Time for asynchronous V1 transfer and execute (ms): %f\n", ms);
  87.   printf("  max error: %e\n", maxError(a, n));

  88.   // asynchronous version 2:
  89.   // loop over copy, loop over kernel, loop over copy
  90.   memset(a, 0, bytes);
  91.   checkCuda( cudaEventRecord(startEvent,0) );
  92.   for (int i = 0; i < nStreams; ++i)
  93.   {
  94.     int offset = i * streamSize;
  95.     checkCuda( cudaMemcpyAsync(&d_a[offset], &a[offset],
  96.                                streamBytes, cudaMemcpyHostToDevice,
  97.                                stream[i]) );
  98.   }
  99.   for (int i = 0; i < nStreams; ++i)
  100.   {
  101.     int offset = i * streamSize;
  102.     kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
  103.   }
  104.   for (int i = 0; i < nStreams; ++i)
  105.   {
  106.     int offset = i * streamSize;
  107.     checkCuda( cudaMemcpyAsync(&a[offset], &d_a[offset],
  108.                                streamBytes, cudaMemcpyDeviceToHost,
  109.                                stream[i]) );
  110.   }
  111.   checkCuda( cudaEventRecord(stopEvent, 0) );
  112.   checkCuda( cudaEventSynchronize(stopEvent) );
  113.   checkCuda( cudaEventElapsedTime(&ms, startEvent, stopEvent) );
  114.   printf("Time for asynchronous V2 transfer and execute (ms): %f\n", ms);
  115.   printf("  max error: %e\n", maxError(a, n));

  116.   // cleanup
  117.   checkCuda( cudaEventDestroy(startEvent) );
  118.   checkCuda( cudaEventDestroy(stopEvent) );
  119.   checkCuda( cudaEventDestroy(dummyEvent) );
  120.   for (int i = 0; i < nStreams; ++i)
  121.     checkCuda( cudaStreamDestroy(stream[i]) );
  122.   cudaFree(d_a);
  123.   cudaFreeHost(a);

  124.   return 0;
  125. }
复制代码


问题2:
    根据问题1中的代码内容,本来是想学习一下stream的并行执行的,结果发现我的机器kernel执行的时候有overlap,但是数据的copy却不能与kernel并发,不知道是什么原因造成的。
一下是我用sample工程执行的profiler图


问题比较多,原本我认为是因为我的工程属性造成数据传输无法并行,所以才使用sample中的工程来执行的,结果发现sample中的工程也无法并发传输,然后又发现程序运行的效率差距很大,其中两个程序都没有使用fast-math选项,结果一个问题编程两个了。卡的属性应当是支持并发传输的,deviceQuery.exe的执行结果中 Concurrent copy and kernel execution:  Yes with 1 copy engine(s)。麻烦论坛里的高手帮忙解惑,谢谢了!

本帖子中包含更多资源

您需要 登录 才可以下载或查看,没有帐号?立即注册

x
使用道具 举报 回复
发表于 2013-9-10 15:19:31
楼主您好:

(1)不同的项目配置导致不同的执行时间,是相当正常的一件事
(例如是否优化,编译参数等因素都可能会极大的影响性能。您的自建项目可能没有使用了例子的编译参数。甚至还是默认的debug配置,那么性能差出几倍,很正常)

(2)关于overlap的时候,普通的geforce卡自CUDA 5.0(的配套驱动版本或更高起),将不再提供更好的overlap支持,因为表现出很差的计算和传输overlap甚至完全无法overlap, 相当正常。
(您可以使用Tesla卡,如果您需要此特性)。
(您也可以可以回退到CUDA 5.0配套的显卡驱动的更低版本号,以便继续在geforce卡上使用多个流并行,但推荐购买Tesla卡)

感谢您的来访。
使用道具 举报 回复 支持 反对
发表于 2013-9-10 15:22:58
稍微补充一下横扫斑竹:

LZ提供的两个截图中,程序的收敛精度(max error)不尽相同,并且收敛精度较低的一组使用时间短。这也可能是您用时不同的原因之一。

使用道具 举报 回复 支持 反对
lzhbch 来自手机
4#
发表于 2013-9-10 15:36:53
感谢两位版主的解答,想再多问一句,关于项目配置的问题所有选项都可以在工程属性中修改,因为当初也怀疑过这个问题,所以仔细对比了工程属性中的参数设置,没有发现什么不同,对工程的参数设置版主有什么好的建议吗?对ice版主的建议,我使用fast math参数后还是sample稍微快些,这时两者的计算精度就一致了。
使用道具 举报 回复 支持 反对
发表于 2013-9-10 15:40:15
lzhbch 发表于 2013-9-10 15:36
感谢两位版主的解答,想再多问一句,关于项目配置的问题所有选项都可以在工程属性中修改,因为当初也怀疑过 ...

那不可能,楼主你至少需要改成release编译。

在论坛上这种类似问题,性能差上个5-10倍,基本99%都是用的debug下编译的。
(您想想,您不过是有复制+计算,前者是固定的DMA engine, 性能必然一样。
而且根据你的图看,复制只占据小部分时间,那么自然是kernel执行时间差出个5倍。
而kernel能执行差出5倍,必然是前文说的问题。你觉得呢?)
使用道具 举报 回复 支持 反对
lzhbch 来自手机
6#
发表于 2013-9-10 17:07:47
横扫千军 发表于 2013-9-10 15:40
那不可能,楼主你至少需要改成release编译。

在论坛上这种类似问题,性能差上个5-10倍,基本99%都是用的 ...

谢谢横扫版主,本来工程属性问题应该要自己研究的,我再仔细检查一遍参数吧,再次感谢。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册