时间测试程序

CUDA将GPU称为设备，将CPU称为主机，一般的计算方式是将数据从内存拷贝进GPU内存(显存)，通过GPU计算再拷回内存中。

下面的代码是一个经典的通过GPU进行的向量加法运算

1. #include<cuda_runtime.h>  
2. #include<windows.h>  
3. #include<iostream>  
4. using namespace std;  
5. const int nMax = 30000;  
6. __global__ void addKernel(float *aaa,float *bbb, float *ccc)  
7. {  
8.     //int i = blockIdx.x;  
9.     int i = threadIdx.x + blockIdx.x*blockDim.x;  
10.     ccc[i] = 0;  
11.     if (i < nMax)for (int j = 0; j < 500; j++)ccc[i] += aaa[i] * bbb[i];  
12. }  
13. void add(float *a, float *b,float *c,int i){  
14.     for (int j = 0; j<500; j++) c[i] += a[i] * b[i];  
15. }  
16. int main(){  
17.     float a[nMax], b[nMax], c[nMax];  
18.     float *devA, *devB, *devC;  
19.     clock_t startT, endT;  
20.     for (int i = 0; i < nMax; i++){  
21.         a[i] = i*1.010923;  
22.         b[i] = 2.13*i;  
23.     }  
24.     startT = clock();  
25.     cudaMalloc((void**)&devA, nMax*sizeof(float));  
26.     cudaMalloc((void**)&devB, nMax*sizeof(float));  
27.     cudaMalloc((void**)&devC, nMax*sizeof(float));  
28.     endT = clock();  
29.     cout << "分配设备空间耗时 " << endT - startT << "ms"<<endl;  
30.   
31.   
32.     startT = clock();  
33.     cudaMemcpy(devA, a,nMax*sizeof(float),cudaMemcpyHostToDevice);  
34.     cudaMemcpy(devB, b, nMax*sizeof(float), cudaMemcpyHostToDevice);  
35.     endT = clock();  
36.     cout << "数据从主机写入设备耗时 " << endT - startT << "ms" << endl;  
37.   
38.     startT = clock();  
39.   
40.     cudaEvent_t start1;  
41.     cudaEventCreate(&start1);  
42.     cudaEvent_t stop1;  
43.     cudaEventCreate(&stop1);  
44.     cudaEventRecord(start1, NULL);  
45.   
46.     addKernel<<<60,501>>>(devA, devB, devC);  
47.   
48.     cudaEventRecord(stop1, NULL);  
49.     cudaEventSynchronize(stop1);  
50.     float msecTotal1 = 0.0f;  
51.     cudaEventElapsedTime(&msecTotal1, start1, stop1);  
52.     //cout << msecTotal1 << "ddd" << endl;  
53.     endT = clock();  
54.     cout << "GPU计算耗时 " << msecTotal1 << "ms" << endl;  
55.   
56.     startT = clock();  
57.     cudaMemcpy(c, devC, nMax*sizeof(float), cudaMemcpyDeviceToHost);  
58.     endT = clock();  
59.     cout << "数据从设备写入主机耗时 " << endT - startT << "ms" << endl;  
60.   
61.     cout <<"GPU计算结果 "<< c[nMax - 1] << endl;  
62.     for (int i = 0; i < nMax; i++){  
63.         a[i] = i*1.010923;  
64.         b[i] = 2.13*i;  
65.         c[i] = 0;  
66.     }  
67.     startT = clock();  
68.     for (int i = 0; i < nMax; i++){  
69.         add(a, b, c, i);  
70.     }  
71.     endT = clock();  
72.     cout << "CPU计算耗时 " << endT - startT << "ms" << endl;  
73.     cout << "CPU计算结果 " << c[nMax - 1] << endl;  
74.   
75.         //释放在设备上分配的空间  
76.     cudaFree(devA);  
77.     cudaFree(devB);  
78.     cudaFree(devC);  
79.     cin >> a[0];  
80.     return 0;  
81. }  

82. 上面的代码中使用了一些通用模式

    1，调用cudaMalloc();这个函数在设备（GPU）上分配内存。一般来说为了避免内存泄漏计算完成之后需要通过cudaFree来释放内存空间。

    2，cudaMemcpy(devA, a,nMax*sizeof(float),cudaMemcpyHostToDevice);这个函数用来处理主机和设备之间的数据拷贝。最后的参数cudaMemcpyHostToDevice代码是从主机拷贝去设备，如果需要从设备拷贝数据到主机需要将这个参数改为cudaMemcpyDeviceToHost。

    3，在定义函数的时候在前面加上__global__，则这个函数就是一个在主机上调用，在设备上运行的函数。在上面的代码里，调用__global__函数代码是

    addKernel<<<60,1>>>(devA, devB, devC);

    这里的<<<60,501>>>的意思是，调用函数的时候，开出60个线程格，每个线程格包含501个线程。在global函数中通过代码int i = threadIdx.x + blockIdx.x*blockDim.x;得到当前线程是第几个线程。

    在调用global函数的时候，我们可以通过dim3类型变量修改调用函数的方式

    例如dim3 grid(10,10);addKernel<<<grid,501>>>(devA, devB, devC);这样就可以把按照一维排列的线程块改为在二维空间内排布。函数内可以通过一下代码得到当前线程的标号

    int x = blockIdx.x;

    int y = blockIdx.y;

    int threadId = x + y *gridDim.x;

    CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：

          1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。

          2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。

          3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。拿上图来说，注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4, blockDim.y = 3。

          4. threadIdx：线程索引，前面章节已经详细探讨过了，这里不再赘述。