CUDA C++ PROGRAMMING GUIDE CH2 PROGRAMMING MODEL

这一章主要介绍 CUDA 编程模型的主要概念，详细描述在第三章：programming interface 里。

Kernels

CUDA C++ 通过 kernel 的概念来对 C++ 进行扩展，其特点是：调用时会在 N 个不同的 CUDA 线程上并行执行 N 次。

kernel 定义的时候用 __global__ 来指定，执行时的 CUDA 线程数通过 <<<...>>> 来进行配置（execution configuration syntax），每一个执行 kernel 的线程都会分配一个唯一的线程 ID，这个 ID 可以通过内置变量来进行访问。

以下例子用变量 threadIdx 来把两个向量 A,B 进行相加，结果存到 C 中：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
	int i = threadIdx.x;
	C[i] = A[i] + B[i];
}
int main()
{
	...
	// Kernel invocation with N threads
	VecAdd<<<1, N>>>(A, B, C);
	...
}

Thread Hierarchy

$\text{threadIdx}$ 是一个三元向量，因此可以构成一维二维三维的线程块，叫做 thread block ，对于二维块大小为 $(D_x,D_y)$ 来说，$(x,y)$ 对应 $(x+y{D}_x)$，三维块大小为 $(D_x,D_y,D_y)$ 来说，$(x,y,z)$ 对应 $(x+y{D}_x+z{D}_x{D}_y)$

两矩阵相加的代码如下：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
					   float C[N][N])
{
	int i = threadIdx.x;
	int j = threadIdx.y;
	C[i][j] = A[i][j] + B[i][j];
}
int main()
{
	...
	// Kernel invocation with one block of N*N*1 threads
	int numBlocks = 1;
	dim3 threadsPerBlock(N, N);
	MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
	...
}

每个 block 上的线程数是有限制的，现在的 GPU 一个 block 上有 1024 个线程。

但是一个 kernel 可以被多个大小相同的 thread block 执行，因此 threads 的总数等于每个 block 的 threads 数乘上 blocks 的数量。

<<<numBlocks, threadsPerBlock>>> 中每个 grid 的 block 数（numBlocks）以及每个 block 的 thread 数（threadsPerBlock）的类型可以是 int 或 dim3 。

grid 中的每一个 block 可以通过 blockIdx 变量来获取，thread block 的维度可以通过 blockDim 变量来获取。

使用多个 block 的话，之前的 MatAdd() 的例子可以改成：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
					   float C[N][N])
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int j = blockIdx.y * blockDim.y + threadIdx.y;
	if (i< N && j < N)
		C[i][j] = A[i][j] + B[i][j];
}
int main()
{
	...
	// Kernel invocation
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
	MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
	...
}

在一个 block 中的 thread 可以通过 shared memory 进行写作，通过调用 __syncthreads() 来设置 kernel 中的同步点。在此处，block 中所有的 threads 都必须等待直到任一 thread 被允许执行。

For efficient cooperation, the shared memory is expected to be a low-latency memory near each processor core (much like an L1 cache) and __syncthreads() is expected to be lightweight.

Memory Hierarchy

CUDA 的线程在执行时可以从多个内存空间中进行数据的访问。每个线程有自己的私有本地内存，每一个 thread block 又有共享内存，然后所有的线程都有一个全局的共享内存。

另外还有两个只读的内存空间，可以被所有的线程访问：the constant and texture memory spaces ，global 、constant、texture 内存空间有自己独特的用处，texture 内存空间还对一些特定的数据提供了不同的寻址模式以及 data filtering

The global, constant, and texture memory spaces are persistent across kernel launches by the same application.

Heterogeneous Programming

异构编程

CUDA 编程模型假设 CUDA 线程在一个物理分隔的 device 上进行执行，host 上运行 C++ 程序。

该模型还假设 device 和 device 都维护各自的内存空间，分别叫做 host memory 和 device memory 。因此，一个程序管理的 global, constant, texture memory space 可以通过调用 CUDA runtime 来使得对 kernel 可见，这包括 device memory 的分配与回收以及 host 和 device memory 之间的数据传递。

unified memory 提供了 managed memory 来拼接 host 和 device 内存空间，这对系统中的 CPU 和 GPU 都是可见的，并且被视为一个有着相同地址空间的整体。This capability enables oversubscription of device memory and can greatly simplify the task of porting applications by eliminating the need to explicitly mirror data on host and device

Compute Capability

计算力用 X.Y 表示，X 为 major revision number ，相同的 X 的 GPU 其核心架构是相同的，比如 X=7，架构为 Volta （伏打），6 为 Pascal ，5 为 Maxwell，3 为 Kepler，2 为 Fermi，1 为 Tesla

Y 为 minor revision number ，一般是在核心架构上的改进，Turing 是 7.5，其为基于 Volta 架构的改进