CUDA C++ Programming Guide——编程模型

原创

mb62de8abf75c00 2022-07-26 10:24:53 ©著作权

©著作权归作者所有：来自51CTO博客作者mb62de8abf75c00的原创作品，请联系作者获取转载授权，否则将追究法律责任

本章通过概述CUDA编程模型在C ++中的使用方式，介绍了其主要概念。编程接口中给出了CUDA C ++的广泛描述。本章和下章使用的向量加法示例的完整代码可以在vectorAdd CUDA示例中找到。

kernels

CUDA C ++通过允许程序员定义称为内核的C ++函数来扩展C ++，这些函数在被调用时由N个不同的CUDA线程并行执行N次，而不是像常规C ++函数那样仅执行一次。使用__global__声明说明符定义内核，并使用<<< … >>>执行配置语法指定为给定内核调用执行该内核的CUDA线程数（请参阅C ++语言扩展）。每个执行内核的线程都有一个唯一的线程ID，可通过内置变量在内核内访问该ID。
作为说明，以下示例代码使用内置变量threadIdx将两个大小为N的向量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

为了方便起见，threadIdx是一个三分量向量，因此可以使用一维，二维或三维线程索引来标识线程，从而形成一个一维，二维或三维块。线程，称为线程块。这提供了一种自然的方法来调用跨域中的元素（例如向量，矩阵或体积）的计算。
线程的索引及其线程ID以直接的方式相互关联：对于一维块，它们是相同的；对于大小为（ $CUDA C++ Programming Guide——编程模型_存储空间$ ）的二维块，索引为（x，y）的线程的线程ID为 $CUDA C++ Programming Guide——编程模型_存储空间_02$ ；对于大小为（ $CUDA C++ Programming Guide——编程模型_共享内存_03$ ）的三维块，索引为（x，y，z）的线程的线程ID为（ $CUDA C++ Programming Guide——编程模型_cuda_04$ ）。
例如，以下代码将两个大小为NxN的矩阵A和B相加，并将结果存储到矩阵C中：

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

每个块的线程数是有限制的，因为块的所有线程都应驻留在同一处理器内核上，并且必须共享该内核的有限内存资源。在当前的GPU上，一个线程块最多可以包含1024个线程。但是，内核可以由多个形状相同的线程块执行，因此线程的总数等于每个块的线程数乘以块数。如图4所示，将块组织成一维，二维或三维的线程块网格。网格中的线程块数通常由所处理数据的大小决定，通常超过系统中的处理器数量。

CUDA C++ Programming Guide——编程模型_cuda_05

CUDA C++ Programming Guide——编程模型_存储空间_06

在<<< … >>>语法中指定的每个块的线程数和每个网格的块数可以是int或dim3类型。可以像上面的示例一样指定二维块或网格。可以通过内置的blockIdx变量在内核中访问的一维，二维或三维唯一索引来标识网格内的每个块。线程块的尺寸可通过内置的blockDim变量在内核中访问。扩展前面的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);
    ...
}

通常选择16x16（256个线程）的线程块大小，尽管在这种情况下是任意的。与以前一样，创建的网格具有足够的块，以使每个矩阵元素具有一个线程。为简单起见，此示例假定每个维度中每个网格的线程数可以被该维度中每个块的线程数平均除尽，尽管并非必须如此。线程块需要独立执行：必须能够以任何顺序（并行或串行）执行它们。这种独立性要求允许线程块在任意数量的内核之间以任意顺序进行调度，如图3所示，从而使程序员可以编写随内核数量扩展的代码。块中的线程可以通过一些共享的内存共享数据并通过同步其执行以协调内存访问来进行协作。更准确地说，可以通过调用__syncthreads内在函数来指定内核中的同步点。 __syncthreads充当屏障，在该屏障中，块中的所有线程必须等待，然后才能继续执行任何线程。共享内存给出了使用共享内存的示例。除了__syncthreads之外，合作组API还提供了一组丰富的线程同步原语。为了进行有效的协作，共享内存应该是每个处理器核心附近的低延迟内存（非常类似于L1缓存），而__syncthreads应该是轻量级的。

Memory Hierarchy

CUDA线程在执行过程中可能会从多个内存空间访问数据，如图5所示。每个线程都有专用的本地内存。 每个线程块具有对该块的所有线程可见的共享内存，并且具有与该块相同的生存期。所有线程都可以访问相同的全局内存。所有线程还可以访问两个附加的只读存储空间：常量存储空间和纹理存储空间。全局，常量和纹理内存空间针对不同的内存使用进行了优化（请参阅设备内存访问）。纹理存储器还为某些特定的数据格式提供了不同的寻址模式以及数据过滤（请参见纹理和表面存储器）。全局，常量和纹理存储空间在同一应用程序的内核启动之间是持久的。

线程 --> 专用本地内存、常量存储空间、纹理存储空间

线程块 --> 对该块的所有线程可见的共享内存

网格 --> 全局内存

CUDA C++ Programming Guide——编程模型_cuda_07

Heterogeneous Programming

如图6所示，CUDA编程模型假定CUDA线程在物理上独立的设备上执行，该设备充当运行C ++程序的主机的协处理器。例如，当内核在GPU上执行而其余C ++程序在CPU上执行时，就是这种情况。CUDA编程模型还假定主机和设备都在DRAM中维护自己的独立内存空间，分别称为主机内存和设备内存。因此，程序通过调用CUDA运行时（在编程接口中介绍）来管理内核可见的全局，常量和纹理存储空间。这包括设备内存的分配和释放以及主机与设备内存之间的数据传输。

统一内存提供托管内存从而桥接主机和设备内存空间。系统中的所有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.

CUDA C++ Programming Guide——编程模型_二维_08

Compute Capability

设备的计算能力由版本号表示，有时也称为“ SM版本”。此版本号标识GPU硬件支持的功能，并由应用程序在运行时用于确定当前GPU上可用的硬件功能和/或指令。计算能力包括主要修订号X和次要修订号Y，并由X.Y表示。具有相同主要版本号的设备具有相同的核心体系结构。主要修订号对于基于NVIDIA Ampere GPU架构的设备为8，对于基于Volta架构的设备为7，对于基于Pascal架构的设备为6，对于基于Maxwell架构的设备为5，对于基于开普勒的设备为3架构，2个用于基于Fermi架构的设备，1个用于基于Tesla架构的设备。
次修订号对应于核心体系结构的增量改进，可能包括新功能。Turing是计算能力为7.5的设备的体系结构，并且是基于Volta体系结构的增量更新。启用CUDA的GPU列出了所有启用CUDA的设备及其计算能力。计算能力给出了每种计算能力的技术规格。
注意：请勿将特定GPU的计算能力版本与CUDA版本（例如CUDA 7.5，CUDA 8，CUDA 9）混淆，CUDA版本是CUDA软件平台的版本。应用程序开发人员使用CUDA平台来创建可在许多代GPU架构上运行的应用程序，包括尚未发明的未来GPU架构。尽管新版本的CUDA平台通常通过支持该架构的计算功能版本来添加对新GPU架构的本机支持，但新版本的CUDA平台通常还包括独立于硬件生成的软件功能。
从CUDA 7.0和CUDA 9.0开始，不再支持Tesla和Fermi架构。

专有变量解读：
定位网格中的线程块：通过内置的blockIdx变量在内核中访问的一维，二维或三维唯一索引来标识网格内的每个块
获取线程块的尺寸：线程块的尺寸可通过内置的blockDim变量在内核中访问
定位线程块中的线程：threadIdx是一个三分量向量，可以使用一维，二维或三维线程索引来标识线程，从而形成一个一维，二维或三维线程块