CUDA从入门到放弃（四）：CUDA 编程模式 CUDA Programming Model

CUDA从入门到放弃（四）：CUDA 编程模式 CUDA Programming Model

1 Kernels

CUDA C++ 扩展了 C++，允许定义名为内核的函数，这些函数可以被不同的 CUDA 线程并行执行多次，而不是像普通 C++ 函数那样只执行一次。内核通过 global 声明符定义，执行内核的 CUDA 线程数通过特殊的执行配置语法指定。每个执行内核的线程都有一个唯一的线程 ID，可在内核内部通过内置变量获取。

// 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);
    ...
}

2 Thread Hierarchy 线程结构

为了方便起见，threadIdx 是一个三分量向量，因此可以使用一维、二维或三维线程索引来识别线程，形成一个一维、二维或三维的线程块，称为线程块。

每个线程块内的线程数有上限，因为所有线程需共享同一流式多处理器核心的有限内存资源。目前GPU上，线程块最多包含1024个线程。 尽管如此，一个内核可以由多个相同形状的线程块执行，总线程数等于每个块的线程数乘以块的数量。

线程块组织成一维、二维或三维网格，由数据大小决定，通常超出系统处理器数量。在执行配置语法中指定的线程数和块数可以是int或dim3类型。每个块通过内置的blockIdx变量在内核中按唯一索引识别，线程块维度通过内置的blockDim变量获取。

∕∕ 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);
	...
}

线程块需能独立执行，可任意顺序并行或串行执行。这允许在多个核心上灵活调度线程块，使代码能随核心数扩展。

块内线程通过共享内存和执行同步来协作，使用__syncthreads()函数设置同步点，作为所有线程必须等待的屏障。共享内存应为低延迟内存，类似于L1缓存，且__syncthreads()操作应轻量级。

2-1 Thread Block Clusters 线程块集群

NVIDIA计算能力9.0引入了一个新的层级结构——线程块集群，由线程块组成。线程块集群中的线程块保证在GPU的处理集群上共同调度。集群可以是一维、二维或三维结构，用户可自定义集群中的线程块数量，CUDA中推荐的最大集群大小为8个线程块。对于小于8个多处理器的GPU硬件或MIG配置，最大集群大小会相应减少。可以通过cudaOccupancyMaxPotentialClusterSize API查询特定架构支持的集群大小。

线程块集群可以在内核中通过使用编译时内核属性 cluster_dims(X,Y,Z) 或者使用CUDA内核启动API cudaLaunchKernelEx 来启用。

2-1-1 编译时内核属性启动集群

使用内核属性的集群大小在编译时固定，然后可以使用传统的 <<< , >>> 语法启动内核。如果一个内核使用了编译时集群大小，那么在启动内核时不能修改集群大小。

∕∕ Kernel definition
∕∕ Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{ }
int main()
{
	float *input, *output;
	∕∕ Kernel invocation with compile time cluster size
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);
	∕∕ The grid dimension is not affected by cluster launch, and is still enumerated
	∕∕ using number of blocks.
	∕∕ The grid dimension must be a multiple of cluster size.
	cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

2-1-2 cudaLaunchKernelEx 来启动内核

使用CUDA内核启动API cudaLaunchKernelEx 来启动内核

∕∕ Kernel definition
∕∕ No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{ }
int main()
{
	float *input, *output;
	dim3 threadsPerBlock(16, 16);
	dim3 numBlocks(N ∕ threadsPerBlock.x, N ∕ threadsPerBlock.y);
	∕∕ Kernel invocation with runtime cluster size
	{
		cudaLaunchConfig_t config = {0};
		∕∕ The grid dimension is not affected by cluster launch, and is still enumerated
		∕∕ using number of blocks.
		∕∕ The grid dimension should be a multiple of cluster size.
		config.gridDim = numBlocks;
		config.blockDim = threadsPerBlock;
		cudaLaunchAttribute attribute[1];
		attribute[0].id = cudaLaunchAttributeClusterDimension;
		attribute[0].val.clusterDim.x = 2; ∕∕ Cluster size in X-dimension
		attribute[0].val.clusterDim.y = 1;
		attribute[0].val.clusterDim.z = 1;
		config.attrs = attribute;
		config.numAttrs = 1;
		cudaLaunchKernelEx(&config, cluster_kernel, input, output);
	}
}

3 Memory Hierarchy 内存结构

CUDA线程在执行过程中可能会访问多个内存空间。每个线程都有私有的本地内存。每个线程块都有共享内存，对块内的所有线程可见，并且与块的生命周期相同。线程块集群中的线程块可以对彼此的共享内存执行读、写和原子操作。所有线程都可以访问相同的全局内存。

4 Heterogeneous Programming

CUDA编程模型假定CUDA线程在物理上独立的设备上执行，作为运行C++程序的主机的协处理器，如GPU和CPU的组合使用。主机和设备各自维护独立的DRAM内存空间。程序通过CUDA运行时管理内核可见的内存空间，包括内存分配、释放和主机与设备间的数据传输。统一内存提供托管内存，实现主机与设备内存空间的统一管理，简化了应用程序的移植过程。

参考资料
1 CUDA编程入门
2 CUDA编程入门极简教程
3 CUDA C++ Programming Guide
4 CUDA C++ Best Practices Guide
5 NVIDIA CUDA初级教程视频
6 CUDA专家手册 [GPU编程权威指南]
7 CUDA并行程序设计：GPU编程指南
8 CUDA C编程权威指南