CUDA刷新器：CUDA编程模型
CUDA Refresher: The CUDA Programming Model
CUDA，CUDA刷新器，并行编程
这是CUDA更新系列的第四篇文章，它的目标是刷新CUDA中的关键概念、工具和初级或中级开发人员的优化。
CUDA编程模型提供了GPU体系结构的抽象，它充当了应用程序与其在GPU硬件上的可能实现之间的桥梁。这篇文章概述了CUDA编程模型的主要概念，概述了它如何在通用编程语言如C/C++中暴露出来。
介绍一下CUDA编程模型中常用的两个关键词：主机和设备。
主机是系统中可用的CPU。与CPU相关联的系统内存称为主机内存。GPU被称为设备，GPU内存也被称为设备内存。
要执行任何CUDA程序，有三个主要步骤：
将输入数据从主机内存复制到设备内存，也称为主机到设备传输。
加载GPU程序并执行，在片上缓存数据以提高性能。
将结果从设备内存复制到主机内存，也称为设备到主机传输。
CUDA内核和线程层次结构
图1显示了CUDA内核是一个在GPU上执行的函数。应用程序的并行部分由K个不同的CUDA线程并行执行k次，而不是像常规C/C++函数那样只进行一次。
Figure 1. The kernel is a function executed on the GPU.
每一个CUDA内核都以一个__global__声明说明符开头。程序员通过使用内置变量为每个线程提供唯一的全局ID。
图2. CUDA内核被细分为块。
一组线程称为CUDA块。CUDA块被分组到一个网格中。内核作为线程块的网格来执行（图2）。
每个CUDA块由一个流式多处理器（SM）执行，不能迁移到GPU中的其他SMs（抢占、调试或CUDA动态并行期间除外）。一个SM可以根据CUDA块所需的资源运行多个并发CUDA块。每个内核在一个设备上执行，CUDA支持一次在一个设备上运行多个内核。图3显示了GPU中可用硬件资源的内核执行和映射。
图3. 在GPU上执行内核。
CUDA为线程和块定义了内置的三维变量。线程使用内置的三维变量threadIdx编制索引。三维索引提供了一种自然的方法来索引向量、矩阵和体积中的元素，并使CUDA编程更容易。类似地，块也使用名为blockIdx的内置三维变量编制索引。
以下是几个值得注意的要点：
CUDA架构限制每个块的线程数（每个块限制1024个线程）。
线程块的维度可以通过内置的blockDim变量在内核中访问。
syncu中的线程可以使用syncu函数同步。使用同步线程时，块中的所有线程都必须等待，然后才能继续。
在<<…>>>语法中指定的每个块的线程数和每个网格的块数可以是int或dim3类型。这些三角括号标记从主机代码到设备代码的调用。它也被称为内核启动。
下面用于添加两个矩阵的CUDA程序显示多维blockIdx和threadIdx以及blockDim等其他变量。在下面的例子中，为了便于索引，选择了一个2D块，每个块有256个线程，x和y方向各有16个线程。使用数据大小除以每个块的大小来计算块的总数。
// Kernel - Adding two matrices MatA and MatB
__global__ void MatAdd(float MatA[N][N], float MatB[N][N],
float MatC[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
MatC[i][j] = MatA[i][j] + MatB[i][j];
}
int main()
{
...
// Matrix addition kernel launch from host code
dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
...
}
Memory hierarchy
支持CUDA的GPU有一个内存层次结构，如图4所示。
图4. gpu中的内存层次结构。
以下内存由GPU架构公开：
这些寄存器对每个线程都是私有的，这意味着分配给线程的寄存器对其他线程不可见。编译器决定寄存器的利用率。
一级/共享内存（SMEM）-每个SM都有一个快速的片上草稿行内存，可用作一级缓存和共享内存。CUDA块中的所有线程都可以共享共享内存，在给定SM上运行的所有CUDA块都可以共享SM提供的物理内存资源。。
只读内存每个SM都有一个指令缓存、常量内存、纹理内存和对内核代码只读的RO缓存。
二级缓存二级缓存在所有SMs中共享，因此每个CUDA块中的每个线程都可以访问该内存。nvidiaa100 GPU已经将二级缓存大小增加到40mb，而v100gpu中只有6mb。
全局内存这是位于GPU中的GPU和DRAM的帧缓冲区大小。
NVIDIA CUDA编译器在优化内存资源方面做得很好，但专家CUDA开发人员可以选择有效地使用这种内存层次结构来优化CUDA程序。
计算能力
GPU的计算能力决定了GPU硬件支持的通用规范和可用特性。此版本号可由应用程序在运行时使用，以确定当前GPU上可用的硬件功能或指令。
每个GPU都有一个版本号，表示为X.Y，其中X包括主要修订号，Y包含次要修订号。小版本号对应于架构的增量改进，可能包括新特性。
有关任何支持CUDA的设备的计算能力的更多信息，请参阅CUDA示例代码设备查询。此示例枚举系统中存在的CUDA设备的属性。
摘要
CUDA编程模型提供了一种异构环境，其中主机代码在CPU上运行C/C++程序，内核在物理上分离的GPU设备上运行。CUDA编程模型还假设主机和设备都保持各自独立的内存空间，分别称为主机内存和设备内存。CUDA代码还通过PCIe总线提供主机和设备内存之间的数据传输。
CUDA还公开了许多内置变量，并提供了多维索引的灵活性，以简化编程。CUDA还管理不同的内存，包括寄存器、共享内存和一级缓存、二级缓存和全局内存。高级开发人员可以有效地使用这些内存来优化CUDA程序。