cuda学习心得--4.检验执行时间

cuda里有cudaEvent_t，通过typedef struct CUevent_st *cudaEvent_t; 。CUevent_st这个结构体并没有看到源代码。也不需要仔细关注。

cudaEventCreate()函数能创建一个cudaEvent_t类型事件，如果需要记录时间，参数2就不需要写，当然我们现在就需要记录时间。接着cudaEventRecord()函数来记录流逝的时间。再调用cudaEventSynchronize同步一下，来确保我们能在cudaEventRecord()这个函数调用来获取准备的值。然后调用cudaEventElapsedTime()函数来获取时间。最后调用cudaEventDestroy()函数来销毁创建的事件。

下面演示一个代码：

 1 #include "cuda_runtime.h"
 2 
 3 #include <iostream>
 4 using namespace std;
 5 
 6 static void HandleError( cudaError_t err,
 7                          const char *file,
 8                          int line ) {
 9     if (err != cudaSuccess) {
10         printf( "%s in %s at line %d\n", cudaGetErrorString( err ),
11                 file, line );
12         exit( EXIT_FAILURE );
13     }
14 }
15 #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
16 
17 
18 #define DIM 2046
19 #define N (DIM * DIM)
20 __global__ void add( int *c, const int* a, const int* b )
21 {
22     int x = blockIdx.x;
23     int y = blockIdx.y;
24     int offset = x + y * gridDim.x;
25     c[offset] = a[offset] + b[offset];
26 }
27 
28 int main()
29 {
30     cudaEvent_t timeStartEvent,timeEndEvent;
31     HANDLE_ERROR ( cudaEventCreate( &timeStartEvent, 0) );
32     HANDLE_ERROR ( cudaEventCreate( &timeEndEvent, 0) );
33     
34     HANDLE_ERROR ( cudaEventRecord( timeStartEvent, 0) );
35 
36     // 测试代码开始
37     int *a, *b, *c;
38     int *dev_a, *dev_b, *dev_c;
39 
40     // 在CPU上分配内存
41     a =  new int[N];
42     b =  new int[N];
43     c =  new int[N];
44     // 在GPU上分配内存
45     HANDLE_ERROR( cudaMalloc( (void**)&dev_a, N * sizeof(int) ) );
46     HANDLE_ERROR( cudaMalloc( (void**)&dev_b, N * sizeof(int) ) );
47     HANDLE_ERROR( cudaMalloc( (void**)&dev_c, N * sizeof(int) ) );
48 
49     // 赋值
50     for (int i=0; i<N; i++) 
51     {
52         a[i] = i;
53         b[i] = 2 * i;
54     }
55 
56     HANDLE_ERROR( cudaMemcpy( dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice ) );
57     HANDLE_ERROR( cudaMemcpy( dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice ) );
58 
59     dim3    grid(DIM,DIM);
60     add<<<grid,1>>>( dev_c,dev_a,dev_b );
61     
62     HANDLE_ERROR( cudaMemcpy( c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost ) );
63 
64     bool success = true;
65     for (int i=0; i<N; i++) 
66     {
67         if ((a[i] + b[i]) != c[i]) 
68         {
69             cout << "Error:  "<<a[i]<<" + "<<b[i]<<" != "<<c[i]<<endl;
70             success = false;
71         }
72     }
73     if (success)    cout << "We did it!\n" ;
74 
75     HANDLE_ERROR( cudaFree( dev_a ) );
76     HANDLE_ERROR( cudaFree( dev_b ) );
77     HANDLE_ERROR( cudaFree( dev_c ) );
78 
79     delete[] a;delete[] b; delete[] c;
80     // 测试代码结束
81     HANDLE_ERROR ( cudaEventRecord( timeEndEvent, 0) );
82     HANDLE_ERROR ( cudaEventSynchronize( timeEndEvent ) );
83     float elapsedTime = 0 ;
84     HANDLE_ERROR ( cudaEventElapsedTime( & elapsedTime, timeStartEvent, timeEndEvent ) );
85 
86     cout << "elapsedTime  " << elapsedTime << " ms. ";
87     HANDLE_ERROR( cudaEventDestroy( timeStartEvent ) );
88     HANDLE_ERROR( cudaEventDestroy( timeEndEvent ) );
89     return 0;
90 }

最终从我的机子里测试出来是运行了309.483ms。当我们把36行到80换成CPU的计算时只需要了0.004736ms。从而可以看出2046*2046的数据量计算，GPU只是处理单纯的的数组加法，并不能发挥优势。

GPU的优势到底在哪里，我们得慢慢找，怎么样的计算，才能让GPU发挥比CPU好。

有兴趣的人，可以慢慢发现时间销毁在哪里。可以用nsight来准确定位到GPU的函数消耗。你也可以调试GPU函数，但需要两张显卡，或者一张显卡上有两个GPU。这里就不重复。