CUDA C++ / 基础编程

核函数

作用

调用核函数的时候，代码会被N个CUDA线程执行N次。

 

修饰符

__global__ 返回值 函数名(){

  ...执行代码

}

 

调用

函数名<<<BlockNumber,ThreadNumber>>>();

BlockNumber是块的个数。

ThreadNumber是每一个块中的线程个数。

 

索引结构

块内线程的索引结构（threadIdx）

在一维，二维，三维的块中，threadIdx分别为一维索引，二维索引，三维索引。

threadIdx.x：一维索引

threadIdx.y：二维索引

threadIdx.z：三维索引

网内块的结构（blockIdx）

在一维，二维，三维的网中，blockIdx分别为一维索引，二维索引，三维索引。

blockIdx.x：一维索引

blockIdx.y：二维索引

blockIdx.z：三维索引

blockDim：块的大小（块内线程总数）

 

向量元素与线程之间映射

一维

blockIdx.x * blockDim.x + threadIdx.x

假设有Item0，Item1，Item2，Item3，Item4，Item5，Item6，Item7八个元素。有两个块block1，block2，每个块4个线程。则分配如下

Item0，Item1，Item2，Item3，分配给block1中的0，1，2，3号线程。

Item4，Item5，Item6，Item7，分配给block2中的0，1，2，3号线程。

则block2中的3号线程处理的元素，根据公式得：1*4+3=7，也就是Item7（索引下标从0开始）

三维

threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x

 

网格大小与工作量匹配

匹配情况

每一个线程对应一个元素

不匹配

元素个数与线程总数不匹配，导致有的线程没有工作可干，运行时会出错。

解决方法

将任务总数N传递到核函数中。

网格内的线程，利用 threadIdx + blockIdx*blockDim检查自己是否超过元素总数N。

 

跨网格工作

通常一个线程不是处理一个元素，而是处理多个元素。

因此每次增长的步长为blockDim.x * gridDim.x（也就是一个网格的所有线程数）

此网格的块数GridDim.x为2，块中的线程数blockDim.x为4，所以一个网格的线程总数为2*4

案例

计算10000个元素，使用256个块，每个块32个线程，一轮可以处理8192个元素。那么在第二轮的时候，只需要判断i增加了后是否超过N即可，超过N的线程不执行。

__global__ void doubleElements(int *a, int N)
{
  int i;
  i = blockIdx.x * blockDim.x + threadIdx.x;
  int stride= blockDim.x * gridDim.x;
  for(;i<N;i+=stride){
   if (i < N)
      {
        a[i] *= 2;
      }
  }

 size_t threads_per_block = 256;
  size_t number_of_blocks = 32;
doubleElements<<<number_of_blocks, threads_per_block>>>(a, N);

 

线程同步

块内同步函数

void __syncthreads();

阻塞块内所有线程，当所有线程到达__syncthreads()才会继续往下执行

 

host端同步函数

cudaDeviceSynchronize();

每次启动核函数后调用该函数，可以保证核函数完成，再继续当前线程。