分析一下Cuda c编程指南

分析一下Cuda c编程指南

CUDA C++编程指南

CUDA 模型和接口的编程指南。

与 12.0 版相比的变化

1. 简介

1.1. 使用 GPU 的好处

图形处理单元 （GPU）1在相似的价格和功率范围内提供比 CPU 高得多的指令吞吐量和内存带宽。许多应用程序利用这些更高的功能在 GPU 上运行比在 CPU 上运行得更快（请参阅 GPU 应用程序）。其他计算设备，如FPGA，也非常节能，但编程灵活性远低于GPU。

GPU 和 CPU 之间存在这种功能差异，因为它们的设计考虑了不同的目标。虽然 CPU 旨在尽可能快地执行一系列操作（称为线程），并且可以并行执行几十个这样的线程，但 GPU 旨在擅长并行执行数千个操作（摊销较慢的单线程性能以实现更大的吞吐量）。

GPU专门用于高度并行计算，因此设计使得更多的晶体管专用于数据处理，而不是数据缓存和流量控制。示意图图 1 显示了 CPU 与 GPU 的芯片资源分布示例。

 GPU将更多的晶体管用于数据处理

将更多的晶体管用于数据处理，例如浮点计算，有利于高度并行的计算;GPU 可以通过计算隐藏内存访问延迟，而不是依靠大数据缓存和复杂的流量控制来避免长内存访问延迟，这两者都在晶体管方面都很昂贵。

通常，应用程序具有并行部分和顺序部分的混合，因此系统设计时混合使用 GPU 和 CPU，以最大限度地提高整体性能。具有高度并行性的应用程序可以利用 GPU 的这种大规模并行特性来实现比 CPU 更高的性能。

1.2. CUDA：®通用并行计算平台和编程模型

2006 年 <> 月，NVIDIA 推出了 CUDA，这是一种通用并行计算平台和编程模型，它利用 NVIDIA GPU 中的并行计算引擎以比 CPU 更有效的方式解决许多复杂的计算问题。®®

CUDA 带有一个软件环境，允许开发人员将C++用作高级编程语言。如图 2 所示，支持其他语言、应用程序编程接口或基于指令的方法，例如 FORTRAN、DirectCompute、OpenACC。

 GPU 计算应用程序。CUDA 旨在支持各种语言和应用程序编程接口。

1.3. 可扩展的编程模型

多核 CPU 和众核 GPU 的出现意味着主流处理器芯片现在是并行系统。挑战在于开发能够透明扩展其并行性的应用软件，以利用越来越多的处理器内核，就像 3D 图形应用程序透明地将其并行性扩展到具有广泛变化内核数量的众核 GPU 一样。

CUDA 并行编程模型旨在克服这一挑战，同时为熟悉标准编程语言（如 C）的程序员保持较低的学习曲线。

它的核心是三个关键抽象——线程组的层次结构、共享内存和屏障同步——它们只是作为一组最小的语言扩展暴露给程序员。

这些抽象提供细粒度数据并行性和线程并行性，嵌套在粗粒度数据并行性和任务并行性中。它们指导程序员将问题划分为粗略的子问题，这些子问题可以通过线程块并行独立解决，并将每个子问题划分为可以由块内的所有线程并行协作解决的更精细的部分。

这种分解通过允许线程在解决每个子问题时进行协作来保持语言表达能力，同时启用自动可伸缩性。实际上，每个线程块都可以以任何顺序、并发或顺序在 GPU 内的任何可用多处理器上调度，以便编译的 CUDA 程序可以在任意数量的多处理器上执行，如图 3 所示，只有运行时系统需要知道物理多处理器计数。

这种可扩展的编程模型允许 GPU 架构通过简单地扩展多处理器和内存分区的数量来跨越广泛的市场范围：从高性能发烧友 GeForce GPU 和专业的 Quadro 和 Tesla 计算产品到各种廉价的主流 GeForce GPU（有关所有支持 CUDA 的 GPU 的列表，请参阅 CUDA 支持的 GPU）。

 自动可扩展性

注意

GPU 是围绕流式多处理器 （SM） 阵列构建的（有关更多详细信息，请参阅硬件实现）。多线程程序被划分为彼此独立执行的线程块，因此具有更多多处理器的 GPU 将在比具有较少多处理器的 GPU 更短的时间内自动执行该程序。

1.4. 文档结构

本文档分为以下部分：

1

图形限定符来自这样一个事实，即二十年前最初创建 GPU 时，它被设计为加速图形渲染的专用处理器。在市场对实时、高清、3D 图形的永不满足需求的推动下，它已经发展成为一种通用处理器，用于处理更多的工作负载，而不仅仅是图形渲染。

2. 编程模型

本章通过概述如何在 C++ 中公开 CUDA 编程模型的主要概念来介绍它们。

CUDA C++的详尽描述在编程接口中给出。

本章和下一章中使用的矢量加法示例的完整代码可以在 vectorAdd CUDA 示例中找到。

2.1. 内核

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

在这里，执行VecAdd()的 N 个线程中的每一个都执行一个成对加法。

2.2. 线程层次结构

为方便起见，threadIdx是一个 3 分量向量，以便可以使用一维、二维或三维线程索引来识别线程，形成一维、二维或三维线程块，称为线程块。这提供了一种自然的方式来跨域（如向量、矩阵或体积）中的元素调用计算。

线程的索引及其线程 ID 以简单的方式相互关联：对于一维块，它们是相同的;对于大小为 （Dx， Dy） 的二维块，索引 （x， y） 的线程的线程 ID 为 （x + y Dx）;对于大小为 （Dx， Dy， Dz） 的三维块，索引 （x， y， z） 的线程的线程 ID 为 （x + y Dx + z Dx Dy）。

例如，以下代码添加两个大小为 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 所示。网格中的线程块数通常由正在处理的数据的大小决定，通常超过系统中的处理器数。

 线程块网格

<<<...>>>语法中指定的每个块的线程数和每个网格的块数可以是 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 还提供了一组丰富的线程同步原语。

为了高效协作，共享内存应是每个处理器内核附近的低延迟内存（非常类似于一级缓存），并且__syncthreads()预计是轻量级的。

2.2.1. 线程块簇

随着 NVIDIA 计算能力 9.0 的引入，CUDA 编程模型引入了一个可选的层次结构级别，称为线程块集群，由线程块组成。与线程块中的线程保证在流式多处理器上共同调度的方式类似，集群中的线程块也保证在 GPU 中的 GPU 处理集群 （GPC） 上共同调度。

与线程块类似，集群也被组织成一维、二维或三维，如图 5 所示。集群中的线程块数量可以由用户定义，并且集群中最多支持 8 个线程块作为 CUDA 中的可移植集群大小。 请注意，在太小而无法支持 8 个多处理器的 GPU 硬件或 MIG 配置上，最大集群大小将相应减小。识别这些较小的配置，以及支持线程块簇大小超过 8 的较大配置，是特定于体系结构的，可以使用cudaOccupancyMaxPotentialClusterSize API 进行查询。

 线程块簇的网格

注意

在使用集群支持启动的内核中，出于兼容性目的，gridDim 变量仍然以线程块的数量表示大小。可以使用集群组 API 找到集群中块的排名。

可以使用编译器时间内核属性__cluster_dims__(X,Y,Z)在内核中启用线程块集群，也可以使用 CUDA 内核启动 API cudaLaunchKernelEx。下面的示例显示了如何使用编译器时间内核属性启动集群。使用内核属性的簇大小在编译时是固定的，然后可以使用经典的<<< , >>>。如果内核使用编译时簇大小，则在启动内核时无法修改簇大小。

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

线程块簇大小也可以在运行时设置，并且可以使用 CUDA 内核启动 API cudaLaunchKernelEx启动内核。下面的代码示例显示了如何使用可扩展 API 启动群集内核。

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

在具有计算能力 9.0 的 GPU 中，保证群集中的所有线程块在单个 GPU 处理群集 （GPC） 上共同调度，并允许群集中的线程块使用群集组 API cluster.sync() 执行硬件支持的同步。集群组还提供成员函数，分别根据线程数或块数num_threads()和num_blocks()API 查询集群组大小。集群组中线程或块的排名可以分别使用dim_threads() 和dim_blocks() API 查询。

属于群集的线程块可以访问分布式共享内存。群集中的线程块能够对分布式共享内存中的任何地址进行读取、写入和执行原子操作。分布式共享内存提供了在分布式共享内存中执行直方图的示例。

2.3. 内存层次结构

CUDA 线程在执行期间可能会从多个内存空间访问数据，如图 6 所示。每个线程都有专用本地内存。每个线程块都具有共享内存，对块的所有线程可见，并且具有与块相同的生存期。线程块群集中的线程块可以对彼此的共享内存执行读取、写入和原子操作。所有线程都可以访问相同的全局内存。

还有两个额外的只读内存空间可供所有线程访问：常量和纹理内存空间。全局、常量和纹理内存空间针对不同的内存使用情况进行了优化（请参阅设备内存访问）。纹理内存还为某些特定数据格式提供不同的寻址模式以及数据过滤（请参阅纹理和表面内存）。

全局、常量和纹理内存空间在同一应用程序的内核启动中是持久的。

 内存层次结构

2.4. 异构编程

如图 7 所示，CUDA 编程模型假设 CUDA 线程在物理上独立的设备上执行，该设备作为运行C++程序的主机的协处理器运行。例如，当内核在 GPU 上执行，而C++程序的其余部分在 CPU 上执行时，就是这种情况。

CUDA 编程模型还假设主机和设备在 DRAM 中维护自己单独的内存空间，分别称为主机内存和设备内存。因此，程序通过调用 CUDA 运行时来管理内核可见的全局、常量和纹理内存空间（如编程接口中所述）。这包括设备内存分配和释放以及主机和设备内存之间的数据传输。

统一内存提供托管内存来桥接主机和设备内存空间。可以从系统中的所有 CPU 和 GPU 访问托管内存，作为具有公共地址空间的单个连贯内存映像。此功能支持设备内存的超额订阅，并且无需在主机和设备上显式镜像数据，从而大大简化了移植应用程序的任务。有关统一内存的简介，请参阅统一内存编程。

 异构编程

注意

串行代码在主机上执行，而并行代码在设备上执行。

 

参考文献链接

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

posted @ 2023-07-04 04:33  吴建明wujianming  阅读(155)  评论(0)  编辑  收藏  举报