2021SC@SDUSC-cuda编程

2021SC@SDUSC

CUDA线程模型

线程是程序执行的最基本单元，CUDA的并行计算就是通过成千上万个线程的并行执行来实现的。
CUDA的线程模型自底向上依次为：

Thread：线程，并行的基本单位
Thread Block：线程块，互相合作的线程组，线程块有如下几个特点：

允许彼此同步
可以通过共享内存快速交换数据
以1维、2维或3维组织

Grid：一组线程块

以1维、2维组织
共享全局内存

Kernel：在GPU上执行的核心程序，这个kernel函数是运行在某个Grid上的。

One kernel <-> One Grid

每一个block和每个thread都有自己的ID，我们通过相应的索引找到相应的线程和线程块。

threadIdx，blockIdx
Block ID: 1D or 2D
Thread ID: 1D, 2D or 3D

另外两个重要概念

• SP：最基本的处理单元，streaming processor，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

• SM：多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。
  

• 每个线程由每个线程处理器（SP）执行

• 线程块由多核处理器（SM）执行

• 一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行

CUDA编程模型

第一个编程要点：

通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑。如下表所示，比如我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，注意__global__函数的返回值必须设置为void。

第二个编程要点：CPU和GPU间的数据传输

GPU内存分配回收内存的函数接口：
• cudaMalloc(): 在设备端分配global memory
• cudaFree(): 释放存储空间
CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：
cudaMemcpy(void *dst, void *src, size_t nbytes,
enum cudaMemcpyKind direction)
enum cudaMemcpyKind:
• cudaMemcpyHostToDevice（CPU到GPU）
• cudaMemcpyDeviceToHost（GPU到CPU）
• cudaMemcpyDeviceToDevice（GPU到GPU）

第三个编程要点：用代码表示线程组织模型

用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

计算线程号

1.使用N个线程块，每一个线程块只有一个线程，即

dim3 dimGrid(N);
dim3 dimBlock(1);

此时的线程号的计算方式就是
threadId = blockIdx.x;
其中threadId的取值范围为0到N-1。对于这种情况，可以将其看作是一个列向量，列向量中的每一行对应一个线程块。列向量中每一行只有1个元素，对应一个线程。
2.使用M×N个线程块，每个线程块1个线程
由于线程块是2维的，故可以看做是一个M*N的2维矩阵，其线程号有两个维度，即：

dim3 dimGrid(M,N);
dim3 dimBlock(1);

其中
blockIdx.x 取值0到M-1
blcokIdx.y 取值0到N-1

这种情况一般用于处理2维数据结构，比如2维图像。每一个像素用一个线程来处理，此时需要线程号来映射图像像素的对应位置，如
pos = blockIdx.y * blcokDim.x + blockIdx.x; //其中gridDim.x等于M
3.使用一个线程块，该线程具有N个线程，即

dim3 dimGrid(1);
dim3 dimBlock(N);

此时线程号的计算方式为
threadId = threadIdx.x;
其中threadId的范围是0到N-1，对于这种情况，可以看做是一个行向量，行向量中的每一个元素的每一个元素对应着一个线程。
4.使用M个线程块，每个线程块内含有N个线程，即

dim3 dimGrid(M);
dim3 dimBlock(N);

这种情况，可以把它想象成二维矩阵，矩阵的行与线程块对应，矩阵的列与线程编号对应，那线程号的计算方式为
threadId = threadIdx.x + blcokIdx*blockDim.x;
上面其实就是把二维的索引空间转换为一维索引空间的过程。
5.使用M×N的二维线程块，每一个线程块具有P×Q个线程，即

dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);

这种情况其实是我们遇到的最多情况，特别适用于处理具有二维数据结构的算法，比如图像处理领域。
其索引有两个维度

threadId.x = blockIdx.xblockDim.x+threadIdx.x; threadId.y = blockIdx.yblockDim.y+threadIdx.y;

上述公式就是把线程和线程块的索引映射为图像像素坐标的计算方法。