CUDA2.3-原理之任意长度的矢量求和与用事件来测量性能

__global__ void add( int *a, int *b, int *c) {
<span style="white-space:pre">	</span> int tid = threadIdx.x + blockIdx.x *blockDim.x;
<span style="white-space:pre">	</span>while (tid < N) {
<span style="white-space:pre">	</span>    c[tid] = a[tid] + b[tid];//your operation
<span style="white-space:pre">	</span>    tid += blockDim.x * gridDim.x;
        }
}

来自《GPU高性能编程CUDA实战》之5.2.1节，这里是采用一维数组（也就是向量）例子，而不是二维矩阵的形式。

 对于任何一个显卡来说，都有硬件的限制，以前的显卡在线程块的数量上有65535的上限，而且每个线程块中的线程数量也有限制，具体的数值可以通过设备属性结构中的maxThreadsPerBlock域的值来查找，该值可以通过函数cudaGetDeviceProperties来得到：

cudaDeviceProp prop;//声明一个结构体来接收显卡属性信息

cudaGetDeviceProperties( &prop, 0 );//调用函数将第0块显卡属性信息提取出来

printf( "Max threads per block:%d\n",prop.maxThreadsPerBlock );//显示每个块上线程数量的上限

 

这里介绍两个新的内置变量 blockDim和gridDim（这两个变量都有.x;.y;.z三个域表示三个不同维度）；前者对于所有的线程块来说，都是一个常数，表示一个线程块中每一维的线程数量；后者对于所有的网格来说，都是一个常数，表示一个网格中每一维的线程块的数量，不过通常来说是二维的网格，所以.z一般都是1。

可以通过对核函数的调用来进行分配线程，一个较为简单的方法是：

add<<<(N+127)/128,128>>>(dev_a,dev_b,dev_c);

将每个块设置成128（举例而已）个线程数量，而N表示你的矩阵的所需总的线程数，当整除的时候刚刚好，不能整除的时候，会多开一个线程块，并通过下面的方法来进行限制：

If(tid<N)//tid表示当前是第几个线程数

/*your operation*/;

换句话说就是当执行的线程数量超过总的所需的线程数量的时候不会有任何的操作。

不过因为网格中的线程块有限制，而线程块中的线程有限制，也就是说当所需要线程的个数超过65535×128=8388460的时候，核函数会调用失败，所以还是需要更大限度的上限。可以用如下的方法：

上面就是核函数的原理代码，意图通过将并行化过程与硬件的实际执行过程解耦开来（就是写代码的时候可以不需要考虑硬件有几个流处理器，也就是之前说的透明性和程序的可扩展性），减轻CUDA程序员的负担。上面就是两个索引的迭代作为核心原理部分：

int tid = threadIdx.x + blockIdx.x *blockDim.x;

 

通过kernel函数的线程索引，这里以块的索引号乘以块的x上的维度的大小（这个块中的线程数量，因为这里是以一维向量做例子，所以只有x，没有y），那么就是跳过了前面的块，然后加上当前块中的线程索引，那么tid就是所需要计算的线程的索引了；这里使用while（）{}的方法是使用分网格的方式，不过网格本就是分开的，因为kernel的<<<>>>两个参数中，第一个是网格的块的维度gridDim，第二个是块的线程维度blockDim，所以没有网格的维度，也就是说上面的tid是第一个网格的线程索引，然后通过加上while（）{}中的代码：tid += blockDim.x * gridDim.x;来将当前网格维度的线程索引以网格维度乘以网格的块的维度（就是每个线程块中的线程数量乘以网格中线程块的数量）来进行递增，就是跳到另一个网格上去（当然了，这是逻辑上的，不是物理上的），既然这样能够无视网格中的块的限制，那么就可以将kernel的参数进行自定义的限制了：add<<<128,128>>>(dev_a,dev_b,dev_c);（128是自定义的，如果可以最好按照上一节中的存储器 和warp的要求来设置）那么这时候这个一维数组的长度只取决于GPU上存储器的容量了，可以通过下面的方法来验证是否成功了：

// verify that the GPU did the work werequested
bool success = true;/设置标识
  for (int i=0; i<N; i++) {
    if ((a[i] + b[i]) != c[i]) {
       printf( “Error: %d + %d != %d\n”, a[i], b[i], c[i] );//这里的操作是加操作
       success = false;//如果没成功那么将这个标识设置成失败
    }
  }
if (success) printf( "We did it!\n" );

 

用事件来测量性能

在cpu上可以使用cpu或者操作系统的计时器来计算一个任务的执行时间，而先不说这有各种延迟（os的线程调度或者高精度cpu计时器的可用性），而且gpu核函数运行时，主机还是可以异步执行计算，所以没法用以往的方式来测量gpu上某个任务上花费的时间，所以需要使用CUDA的事件API来解决这个问题。

其本质是一个GPU时间戳，只需要两个步骤：创建一个事件，然后记录这个事件：

cudaEvent_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);

cudaEventRecord( start, 0 );

// do some work on the GPU

cudaEventRecord( stop, 0 );

先创建开始和终止的事件，然后调用记录事件函数（这里的第二个参数0 暂时先不说，记住就行，后续在介绍）记录开始的事件，然后进行操作，接着调用记录事件函数记录终止的事件。

但是因为cpu和gpu可以进行异步的函数调用，也就是说当gpu开始执行代码，但是未执行完成的时候，cpu却能够继续执行程序中的下一行代码，从性能的角度看是非常不错的，但是从逻辑角度来看，这样就使得计时工作没法准确的进行下去（中间虽然说是填写do work on gpu的代码，但是这部分难免会有在主机上的代码，比如先在主机上分配内存和数据然后在复制到设备上）。所以还需要将它们进行同步起来，所以需要在后面加上一个函数：cudaEventSynchronize( stop );

，最后的结果如下：

cudaEvent_t start, stop;

cudaEventCreate(&start);

cudaEventCreate(&stop);

cudaEventRecord( start, 0 );

// do some work on the GPU

cudaEventRecord( stop, 0 );

cudaEventSynchronize( stop );

float elapsedTime;

cudaEventElapsedTime( &elapsedTime,start, stop ) ;

printf( "Time to generate: %3.1f ms\n", elapsedTime );

cudaEventDestroy( start );

cudaEventDestroy( stop ) ;

这样会告诉运行时系统阻塞后面的代码，而且当cudaEventSynchronize( stop );函数返回的时候，就知道在stop事件之前的所有gpu工作都已经完成了，可以安全的读取在stop中保持的时间戳，不过记得，因为cuda事件是直接在gpu上实现的，所以不适合用于同时包含设备代码和主机代码的混合代码计时，也就是说，不能依靠这个来通过cuda事件对核函数和设备内存复制之外的代码进行计时，因为结果不准（个人：这个函数的意思应该是让cpu和gpu同步起来，使得这个函数前面的代码都能够执行完成，而且cpu不会跳到后面去执行，但是因为cpu中会有延时或者什么，因为毕竟这里的事件是在gpu上执行的，所以估计这才是不准确的原因吧，不过关于这个函数的更详细的具体意义可以查看cuda自带的文档）。这里的cudaEventElapsedTime()和cudaEventDestory()是两个收尾的函数，前者是一个工具函数用来计算两个事件之间经历的时间，第一个参数是某个浮点变量的地址，其中记录着两个事件之间的时间，单位为毫秒；后个函数就是简单的释放掉事件的函数