常看到有人说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。
测试程序如下,欢迎指正。
- #include "cutil_inline.h"
- #include "stdio.h"
- #include "stdlib.h"
- #define NN 10000000
- #define I (240*64)
- #define K 2000
- #define RAND() ( (RAND_MAX<=32768) ? rand()*rand()+1 : rand() )
- int res1[I], res2[I], link[NN], *d_res, *d_link;
- __global__ void gputest(int *d_link, int *d_res) {
- int inx=blockIdx.x*blockDim.x+threadIdx.x;
- int p=inx;
- for(int k=0; k<K; k++) p=d_link[p];
- d_res[inx]=p;
- }
- int main() {
- double t0, t1, dt_cpu, dt_gpu, result_cpu, result_gpu;
- unsigned int timer1;
- cutilSafeCall(cudaMalloc((void**)&d_res, sizeof(res2)));
- cutilSafeCall(cudaMalloc((void**)&d_link, sizeof(link)));
- srand(41);
- for(int i=0; i<NN; i++) link[i]=i;
- for(int i=0; i<NN; i++) { //初始化
- int j=RAND()%NN;
- int tmp=link[i];
- link[i]=link[j];
- link[j]=tmp;
- }
- cutCreateTimer(&timer1);
- cutStartTimer(timer1);
- t0=cutGetTimerValue(timer1);
- for(int i=0; i<I; i++) { //CPU部分
- int p=i;
- for(int k=0; k<K; k++) p=link[p];
- res1[i]=p;
- }
- t1=cutGetTimerValue(timer1);
- dt_cpu=(t1-t0)/1000.0;
- result_cpu=0; for(int i=0; i<I; i++) result_cpu+=res1[i];
- printf("CPU dt=%.3fs, result=%d\n", dt_cpu, result_cpu);
- printf("CPU有效带宽: %.3fMB/s\n", (1e-6*4*I*K)/dt_cpu);
- cutilSafeCall(cudaMemcpy(d_link, link, sizeof(link), cudaMemcpyHostToDevice));
- t0=cutGetTimerValue(timer1);
- gputest<<<240, 64>>>(d_link, d_res); //GPU部分
- cudaThreadSynchronize();
- t1=cutGetTimerValue(timer1);
- dt_gpu=(t1-t0)/1000.0;
- cutilSafeCall(cudaMemcpy(res2, d_res, sizeof(res2), cudaMemcpyDeviceToHost));
- result_gpu=0; for(int i=0; i<I; i++) result_gpu+=res2[i];
- printf("GPU dt=%.3fs, result=%d\n", dt_gpu, result_gpu);
- printf("GPU有效带宽: %.3fMB/s\n", (1e-6*4*I*K)/dt_gpu);
- printf("Test %s, 加速比=%.1fX\n", (result_cpu==result_gpu) ? "PASSED" : "FAILED", dt_cpu/dt_gpu);
- }
复制代码
CUDA 2010 作品! |