用户
 找回密码
 立即注册
cscuda 该用户已被删除
发表于 2010-5-26 14:15:52
79751
常看到有人说CUDA上面gmem的随机访问性能很差,与CPU不在一个量级。但下面这个测试给出了完全相反的结论。
测试环境为GTX295+CUDA3.0+VS2008, CPU为i7 920。(如果是xeon 5335,CPU成绩还要差一半左右)
这个测试里面所有内存访问都是随机的,并且数据前后相关。测试结果是GPU的速度快几十倍。我觉得当数据量比较大的时候,GPU随机访问还是有明显优势的。

程序运行结果:
CPU dt=1.977s, result=-1125187584
CPU有效带宽: 62.162MB/s
GPU dt=0.053s, result=-1125187584
GPU有效带宽: 2326.469MB/s
Test PASSED, 加速比=37.4X

如果参数NN增大为100000000,加速比可进一步提高到47.6。

测试程序如下,欢迎指正。

  1. #include "cutil_inline.h"
  2. #include "stdio.h"
  3. #include "stdlib.h"
  4. #define NN 10000000
  5. #define I  (240*64)
  6. #define K  2000
  7. #define RAND() ( (RAND_MAX<=32768) ? rand()*rand()+1 : rand() )
  8. int res1[I], res2[I], link[NN], *d_res, *d_link;
  9. __global__ void gputest(int *d_link, int *d_res) {
  10. int inx=blockIdx.x*blockDim.x+threadIdx.x;
  11. int p=inx;
  12. for(int k=0; k<K; k++) p=d_link[p];
  13. d_res[inx]=p;
  14. }
  15. int main() {
  16.     double t0, t1, dt_cpu, dt_gpu, result_cpu, result_gpu;
  17.     unsigned int timer1;
  18. cutilSafeCall(cudaMalloc((void**)&d_res, sizeof(res2)));
  19. cutilSafeCall(cudaMalloc((void**)&d_link, sizeof(link)));
  20. srand(41);
  21. for(int i=0; i<NN; i++) link[i]=i;
  22.     for(int i=0; i<NN; i++) { //初始化
  23.         int j=RAND()%NN;
  24.         int tmp=link[i];
  25.         link[i]=link[j];
  26.         link[j]=tmp;
  27.     }
  28. cutCreateTimer(&timer1);
  29. cutStartTimer(timer1);
  30.     t0=cutGetTimerValue(timer1);
  31. for(int i=0; i<I; i++) { //CPU部分
  32.         int p=i;
  33.         for(int k=0; k<K; k++) p=link[p];
  34.   res1[i]=p;
  35.     }
  36. t1=cutGetTimerValue(timer1);
  37. dt_cpu=(t1-t0)/1000.0;
  38.     result_cpu=0; for(int i=0; i<I; i++) result_cpu+=res1[i];
  39.     printf("CPU dt=%.3fs, result=%d\n", dt_cpu, result_cpu);
  40. printf("CPU有效带宽: %.3fMB/s\n", (1e-6*4*I*K)/dt_cpu);

  41. cutilSafeCall(cudaMemcpy(d_link, link, sizeof(link), cudaMemcpyHostToDevice));
  42. t0=cutGetTimerValue(timer1);
  43. gputest<<<240, 64>>>(d_link, d_res); //GPU部分
  44. cudaThreadSynchronize();
  45. t1=cutGetTimerValue(timer1);
  46. dt_gpu=(t1-t0)/1000.0;
  47. cutilSafeCall(cudaMemcpy(res2, d_res, sizeof(res2), cudaMemcpyDeviceToHost));
  48. result_gpu=0; for(int i=0; i<I; i++) result_gpu+=res2[i];
  49. printf("GPU dt=%.3fs, result=%d\n", dt_gpu, result_gpu);
  50. printf("GPU有效带宽: %.3fMB/s\n", (1e-6*4*I*K)/dt_gpu);
  51. printf("Test %s, 加速比=%.1fX\n", (result_cpu==result_gpu) ? "PASSED" : "FAILED", dt_cpu/dt_gpu);
  52. }
复制代码


CUDA 2010 作品!
使用道具 举报 回复
发表于 2012-10-10 12:39:10
牛,期待更多实验。
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册