CUDA C++ 编程指南学习（待更）

CUDA C++ 编程指南 (nvidia.com)https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html

2. 编程模型

2.1. 内核

CUDA C++ 扩展了 C++，允许程序员定义 C++ 函数，称为内核，当被调用时，N 个不同的 CUDA 线程并行执行 N 次，而不是像常规 C++ 函数那样只执行一次。

内核是使用声明说明符定义的，对于给定的内核调用执行该内核的 CUDA 线程数是使用新的执行配置语法指定的（请参阅 C++ 语言扩展）。每个执行内核的线程都被赋予一个唯一的线程 ID，可以通过内置变量在内核内访问该 ID。__global__<<<...>>>

举例来说，以下示例代码使用内置变量 ，将两个大小为 N 的向量 A 和 B 相加，并将结果存储到向量 C 中：threadIdx

#include <iostream>
#include <cuda_runtime.h>#define N 10 // Vector size// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{int i = threadIdx.x;if (i < N) // Ensure index is within bounds{C[i] = A[i] + B[i];}
}int main()
{// Allocate host memoryfloat *h_A = new float[N];float *h_B = new float[N];float *h_C = new float[N];// Initialize host vectorsfor (int i = 0; i < N; ++i){h_A[i] = i * 1.0f; // Example valuesh_B[i] = i * 2.0f;}// Allocate device memoryfloat *d_A, *d_B, *d_C;cudaMalloc(&d_A, N * sizeof(float));cudaMalloc(&d_B, N * sizeof(float));cudaMalloc(&d_C, N * sizeof(float));// Copy host vectors to devicecudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);// Kernel invocation with N threadsVecAdd<<<1, N>>>(d_A, d_B, d_C);// Copy result from device to hostcudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);// Print resultstd::cout << "Result vector C:" << std::endl;for (int i = 0; i < N; ++i){std::cout << h_C[i] << " ";}std::cout << std::endl;// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memorydelete[] h_A;delete[] h_B;delete[] h_C;return 0;
}

输出：

Result vector C:
0 3 6 9 12 15 18 21 24 27 

2.2. 线程层次结构

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

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

线程的索引和它的线程 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 中：

#include <iostream>
#include <cuda_runtime.h>#define N 3 // Matrix size// Kernel definition
__global__ void MatAdd(float *A, float *B, float *C, int n)
{int i = threadIdx.x;int j = threadIdx.y;if (i < n && j < n) // Ensure index is within bounds{int index = i * n + j; // Flattened index for 2D accessC[index] = A[index] + B[index];}
}int main()
{// Allocate host memoryfloat *h_A = new float[N * N];float *h_B = new float[N * N];float *h_C = new float[N * N];// Initialize host matricesfor (int i = 0; i < N; ++i){for (int j = 0; j < N; ++j){h_A[i * N + j] = static_cast<float>(i + j); // Example initializationh_B[i * N + j] = static_cast<float>(i - j); // Example initialization}}// Allocate device memoryfloat *d_A, *d_B, *d_C;cudaMalloc(&d_A, N * N * sizeof(float));cudaMalloc(&d_B, N * N * sizeof(float));cudaMalloc(&d_C, N * N * sizeof(float));// Copy host matrices to devicecudaMemcpy(d_A, h_A, N * N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, N * N * sizeof(float), cudaMemcpyHostToDevice);// Kernel invocation with one block of N * N threadsdim3 threadsPerBlock(N, N);MatAdd<<<1, threadsPerBlock>>>(d_A, d_B, d_C, N);// Copy result from device to hostcudaMemcpy(h_C, d_C, N * N * sizeof(float), cudaMemcpyDeviceToHost);// Print resultstd::cout << "Result matrix C:" << std::endl;for (int i = 0; i < N; ++i){for (int j = 0; j < N; ++j){std::cout << h_C[i * N + j] << " ";}std::cout << std::endl;}// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memorydelete[] h_A;delete[] h_B;delete[] h_C;return 0;
}

 输出：

Result matrix C:
0 0 0 
2 2 2 
4 4 4 

每个块的线程数是有限制的，因为一个块的所有线程都应该驻留在同一个流式多处理器核心上，并且必须共享该核心的有限内存资源。在当前 GPU 上，一个线程块最多可以包含 1024 个线程。

但是，一个内核可以由多个形状相等的线程块执行，因此线程总数等于每个块的线程数乘以块的数量。

块被组织成一维、二维或三维的螺纹块网格，如图 4 所示。网格中的线程块数量通常由正在处理的数据的大小决定，该大小通常超过系统中的处理器数量。

语法中指定的每个块的线程数和每个网格的块数可以是 或 类型。可以指定二维块或网格，如上例所示。<<<...>>>intdim3

网格中的每个块都可以通过一维、二维或三维唯一索引来识别，该索引可通过内置变量在内核内访问。线程块的维度可以通过内置变量在内核中访问。blockIdxblockDim

扩展上一个示例以处理多个块，代码如下所示。MatAdd()

#include <iostream>
#include <cuda_runtime.h>#define N 32 // Matrix size, must be divisible by threadsPerBlock dimensions// Kernel definition
__global__ void MatAdd(float *A, float *B, float *C, int n)
{int i = blockIdx.x * blockDim.x + threadIdx.x;int j = blockIdx.y * blockDim.y + threadIdx.y;if (i < n && j < n) // Ensure index is within bounds{int index = i * n + j; // Flattened index for 2D accessC[index] = A[index] + B[index];}
}int main()
{// Allocate host memoryfloat *h_A = new float[N * N];float *h_B = new float[N * N];float *h_C = new float[N * N];// Initialize host matricesfor (int i = 0; i < N; ++i){for (int j = 0; j < N; ++j){h_A[i * N + j] = static_cast<float>(i + j); // Example initializationh_B[i * N + j] = static_cast<float>(i - j); // Example initialization}}// Allocate device memoryfloat *d_A, *d_B, *d_C;cudaMalloc(&d_A, N * N * sizeof(float));cudaMalloc(&d_B, N * N * sizeof(float));cudaMalloc(&d_C, N * N * sizeof(float));// Copy host matrices to devicecudaMemcpy(d_A, h_A, N * N * sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, N * N * sizeof(float), cudaMemcpyHostToDevice);// Kernel invocationdim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C, N);// Copy result from device to hostcudaMemcpy(h_C, d_C, N * N * sizeof(float), cudaMemcpyDeviceToHost);// Print resultstd::cout << "Result matrix C:" << std::endl;for (int i = 0; i < N; ++i){for (int j = 0; j < N; ++j){std::cout << h_C[i * N + j] << " ";}std::cout << std::endl;}// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memorydelete[] h_A;delete[] h_B;delete[] h_C;return 0;
}

输出： 

Result matrix C:
0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 
2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 
4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 4 
6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 6 
8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8 
10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 10 
12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 12 
14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 14 
16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 
18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 18 
20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 20 
22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 22 
24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 24 
26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 26 
28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 28 
30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 30 
32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 32 
34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 34 
36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 36 
38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 
40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 40 
42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 42 
44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 44 
46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 46 
48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 48 
50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 50 
52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 52 
54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 54 
56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 56 
58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 58 
60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 60 
62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 62 

线程块大小为 16x16（256 个线程），尽管在这种情况下是任意的，但是一种常见的选择。网格是用足够的块创建的，就像以前一样，每个矩阵元素都有一个线程。为简单起见，此示例假定每个维度中每个网格的线程数可以被该维度中每个块的线程数整除，尽管情况并非如此。

线程块需要独立执行：必须能够以任何顺序执行它们，并行或串联。这种独立性要求允许在任意数量的内核上按任何顺序调度线程块，如图 3 所示，使程序员能够编写随内核数量缩放的代码。

块中的线程可以通过一些共享内存共享数据来协作，并通过同步它们的执行以协调内存访问。更准确地说，可以通过调用内部函数来指定内核中的同步点; 充当一个障碍，块中的所有线程都必须等待该障碍，然后才能允许任何线程继续进行。共享内存给出了使用共享内存的示例。此外，Cooperative Groups API 还提供了一组丰富的线程同步原语。__syncthreads()__syncthreads()__syncthreads()

为了实现高效合作，共享内存应是每个处理器内核附近的低延迟内存（很像 L1 缓存），并且预计是轻量级的。__syncthreads()

2.3. 内存层次结构

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

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

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

2.4. 异构编程 

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

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

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

2.5. 异步SIMT编程模型 

在 CUDA 编程模型中，线程是用于执行计算或内存操作的最低抽象级别。从基于 NVIDIA Ampere GPU 架构的设备开始，CUDA 编程模型通过异步编程模型为内存操作提供加速。异步编程模型定义了异步操作相对于 CUDA 线程的行为。

异步编程模型定义了 CUDA 线程之间同步的异步屏障行为。该模型还解释并定义了 cuda：：memcpy_async 可用于在 GPU 中计算时从全局内存异步移动数据。

 2.5.1. 异步操作

异步操作定义为由 CUDA 线程启动并由另一个线程异步执行的操作，就像其他线程一样。在格式正确的程序中，一个或多个 CUDA 线程与异步操作同步。启动异步操作的 CUDA 线程不需要位于同步线程之间。

此类异步线程（假设线程）始终与启动异步操作的 CUDA 线程相关联。异步操作使用同步对象来同步操作的完成。这样的同步对象可以由用户显式管理（例如，），也可以在库中隐式管理（例如，）。cuda::memcpy_asynccooperative_groups::memcpy_async

同步对象可以是 a 或 。使用 cuda：:p ipeline 的异步屏障和异步数据副本中详细介绍了这些对象。这些同步对象可以在不同的线程作用域中使用。作用域定义了一组线程，这些线程可以使用同步对象与异步操作同步。下表定义了 CUDA C++ 中可用的线程范围以及可以与每个线程同步的线程。cuda::barriercuda::pipeline

 这些线程作用域是作为 CUDA 标准 C++ 库中标准 C++ 的扩展实现的。

2.6. 计算能力 

设备的计算能力由版本号表示，有时也称为其“SM 版本”。此版本号标识 GPU 硬件支持的功能，并在运行时由应用程序用于确定当前 GPU 上可用的硬件功能和/或指令。

计算功能包括一个主要修订号 X 和一个次要修订号 Y，用 X.Y 表示。

具有相同主要修订号的设备具有相同的核心体系结构。对于基于 NVIDIA Hopper GPU 架构的设备，主要修订号为 9，对于基于 NVIDIA Ampere GPU 架构的设备，主要修订号为 8，对于基于 Volta 架构的设备，对于基于 Volta 架构的设备，主要修订号为 6，对于基于 Maxwell 架构的设备，主要修订号为 5，对于基于 Kepler 架构的设备，主要修订号为 3。

次要修订号对应于对核心架构的增量改进，可能包括新功能。

Turing 是计算能力为 7.5 的设备的架构，是基于 Volta 架构的增量更新。

启用 CUDA 的 GPU 列出了所有启用了 CUDA 的设备及其计算能力。计算功能提供了每种计算功能的技术规格。

3. 编程接口 

CUDA C++ 为熟悉 C++ 编程语言的用户提供了一条简单的路径，可以轻松编写供设备执行的程序。

它由 C++ 语言的最小扩展集和一个运行时库组成。

核心语言扩展已在编程模型中引入。它们允许程序员将内核定义为 C++ 函数，并在每次调用函数时使用一些新语法来指定网格和块维度。有关所有扩展的完整说明，请参阅 C++ 语言扩展。包含其中一些扩展的任何源文件都必须按照使用 NVCC 进行编译中所述进行编译。nvcc

运行时是在 CUDA Runtime 中引入的。它提供在主机上执行的 C 和 C++ 函数，用于分配和释放设备内存、在主机内存和设备内存之间传输数据、管理具有多个设备的系统等。可以在 CUDA 参考手册中找到运行时的完整描述。

运行时构建在较低级别的 C API（CUDA 驱动程序 API）之上，应用程序也可以访问该 API。驱动程序 API 通过公开较低级别的概念（例如 CUDA 上下文（设备主机进程的类似物）和 CUDA 模块（设备的动态加载库的类似物）来提供额外的控制级别。大多数应用程序不使用驱动程序 API，因为它们不需要这种额外的控制级别，并且在使用运行时时，上下文和模块管理是隐式的，从而导致代码更简洁。由于运行时可以与驱动程序 API 互操作，因此大多数需要某些驱动程序 API 功能的应用程序可以默认使用运行时 API，并且仅在需要时使用驱动程序 API。驱动程序 API 在驱动程序 API 中介绍，并在参考手册中进行了全面描述。

3.1. 使用NVCC编译 

可以使用称为 PTX 的 CUDA 指令集架构编写内核，PTX 参考手册中对此进行了介绍。但是，使用高级编程语言（如 C++）通常更有效。在这两种情况下，都必须将内核编译为二进制代码才能在设备上执行。nvcc

nvcc是一个编译器驱动程序，可简化编译 C++ 或 PTX 代码的过程：它提供简单熟悉的命令行选项，并通过调用实现不同编译阶段的工具集合来执行它们。本部分概述了工作流和命令选项。完整的说明可以在用户手册中找到。nvccnvcc

3.1.1. 编译工作流程 

3.1.1.1. 离线编译 

编译时使用的源文件可以包含主机代码（即在主机上执行的代码）和设备代码（即在设备上执行的代码）的混合。的基本工作流程包括将设备代码与主机代码分离，然后：nvccnvcc

• 将设备代码编译为汇编形式（PTX 代码）和/或二进制形式（cubin 对象），

• 并通过替换内核中引入的语法（并在执行配置中更详细地描述）来修改主机代码，方法是使用必要的 CUDA 运行时函数调用来从 PTX 代码和/或 cubin 对象加载和启动每个编译的内核。<<<...>>>

修改后的主机代码可以输出为 C++ 代码，然后使用其他工具进行编译，也可以通过在最后一个编译阶段调用主机编译器直接输出为目标代码。nvcc

然后，应用程序可以：

• 链接到已编译的主机代码（这是最常见的情况），

• 或者忽略修改后的主机代码（如果有）并使用 CUDA 驱动程序 API（请参阅驱动程序 API）加载和执行 PTX 代码或 cubin 对象。

 3.1.1.2. 即时编译

应用程序在运行时加载的任何 PTX 代码都会由设备驱动程序进一步编译为二进制代码。这称为实时编译。实时编译会增加应用程序加载时间，但允许应用程序从每个新设备驱动程序附带的任何新编译器改进中受益。这也是应用程序在编译应用程序时不存在的设备上运行的唯一方式，如应用程序兼容性中所述。

当设备驱动程序实时为某些应用程序编译某些 PTX 代码时，它会自动缓存生成的二进制代码的副本，以避免在应用程序的后续调用中重复编译。升级设备驱动程序时，缓存（称为计算缓存）将自动失效，以便应用程序可以从设备驱动程序中内置的新实时编译器的改进中受益。

环境变量可用于控制实时编译，如 CUDA 环境变量中所述

作为用于编译 CUDA C++ 设备代码的替代方法，NVRTC 可用于在运行时将 CUDA C++ 设备代码编译为 PTX。NVRTC 是 CUDA C++ 的运行时编译库;有关更多信息，请参阅 NVRTC 用户指南。nvcc

 3.1.2. 二进制兼容性

 二进制代码是特定于体系结构的。cubin 对象是使用指定目标体系结构的编译器选项生成的：例如，编译 with 会为计算能力为 8.0 的设备生成二进制代码。从一个次要修订版本到下一个修订版本，但不能保证从一个次要修订版本到前一个版本或跨主要修订版本的二进制兼容性。换言之，为计算能力 X.y 生成的立方体对象只会在计算能力为 X.z 的设备上执行，其中 z≥y。-code-code=sm_80

3.1.3. PTX 兼容性 

某些 PTX 指令仅在计算能力较高的设备上受支持。例如，Warp Shuffle Functions 仅在计算能力为 5.0 及以上的设备上受支持。编译器选项指定在将 C++ 编译为 PTX 代码时假定的计算能力。因此，例如，包含 warp shuffle 的代码必须使用 （或更高） 进行编译。-arch-arch=compute_50

为某些特定计算能力生成的 PTX 代码始终可以编译为计算能力更大或相等的二进制代码。请注意，从早期 PTX 版本编译的二进制文件可能无法使用某些硬件功能。例如，从为计算能力 6.0 （Pascal） 生成的 PTX 编译的计算能力 7.0 （Volta） 的二进制目标设备将不会使用 Tensor Core 指令，因为这些指令在 Pascal 上不可用。因此，最终二进制文件的性能可能比使用最新版本的 PTX 生成二进制文件时的性能更差。

为目标架构条件特征编译的 PTX 代码仅在完全相同的物理架构上运行，而不能在其他任何地方运行。Arch 条件 PTX 代码向前和向后不兼容。 使用具有计算能力 9.0 的设备编译的示例代码或仅在具有计算能力 9.0 的设备上运行，并且不向后或向前兼容。sm_90acompute_90a

3.1.4. 应用程序兼容性 

若要在具有特定计算能力的设备上执行代码，应用程序必须加载与此计算功能兼容的二进制代码或 PTX 代码，如二进制兼容性和 PTX 兼容性中所述。具体而言，为了能够在具有更高计算能力的未来架构上执行代码（尚无法生成二进制代码），应用程序必须加载 PTX 代码，这些代码将为这些设备进行实时编译（请参阅实时编译）。

哪些 PTX 和二进制代码嵌入到 CUDA C++ 应用程序中由 和 编译器选项或编译器选项控制，详见用户手册。例如-arch-code-gencodenvcc

nvcc x.cu-gencode arch=compute_50,code=sm_50-gencode arch=compute_60,code=sm_60-gencode arch=compute_70,code=\"compute_70,sm_70\"

嵌入与计算能力 5.0 和 6.0 兼容的二进制代码（第一和第二个选项）以及与计算能力 7.0 兼容的 PTX 和二进制代码（第三个选项）。-gencode-gencode

生成主机代码是为了在运行时自动选择要加载和执行的最合适的代码，在上面的示例中，这些代码将是：

• 具有计算能力 5.0 和 5.2 的设备的 5.0 二进制代码，

• 具有计算能力 6.0 和 6.1 的设备的 6.0 二进制代码，

• 具有计算能力 7.0 和 7.5 的设备的 7.0 二进制代码，

• PTX 代码，在运行时编译为二进制代码，适用于具有计算能力 8.0 和 8.6 的设备。

x.cu可以具有使用变形减少操作的优化代码路径，例如，仅在计算能力为 8.0 和更高功能的设备中受支持。该巨集可用于根据计算能力区分各种代码路径。它仅针对设备代码定义。例如，当编译时，等于 。__CUDA_ARCH__-arch=compute_80__CUDA_ARCH__800

如果使用 或 编译架构条件特性示例，则代码只能在具有计算能力 9.0 的设备上运行。x.cusm_90acompute_90a

使用驱动程序 API 的应用程序必须编译代码以分隔文件，并在运行时显式加载和执行最合适的文件。

Volta 架构引入了独立线程调度，它改变了 GPU 上线程的调度方式。对于依赖于以前体系结构中 SIMT 调度的特定行为的代码，独立线程调度可能会更改参与线程的集合，从而导致错误的结果。为了在实施独立线程调度中详述的纠正措施时帮助迁移，Volta 开发人员可以选择使用编译器选项组合加入 Pascal 的线程调度。-arch=compute_60 -code=sm_70

用户手册列出了 、 和 编译器选项的各种简写。例如，是 的简写（与 相同）。nvcc-arch-code-gencode-arch=sm_70-arch=compute_70 -code=compute_70,sm_70-gencode arch=compute_70,code=\"compute_70,sm_70\"

3.1.5. C++ 兼容性 

 编译器的前端根据 C++ 语法规则处理 CUDA 源文件。主机代码支持完整的 C++。但是，设备代码仅完全支持 C++ 的一个子集，如 C++ 语言支持中所述。

3.1.6. 64位兼容性

64 位版本以 64 位模式编译设备代码（即指针为 64 位）。以 64 位模式编译的设备代码仅支持以 64 位模式编译的主机代码。nvcc 

3.2. CUDA 运行时 

运行时是在库中实现的，该库通过 OR 静态链接到应用程序，或者通过 或 动态链接到应用程序。需要和/或用于动态链接的应用程序通常将它们作为应用程序安装包的一部分包含在内。只有在链接到 CUDA 运行时的同一实例的组件之间传递 CUDA 运行时符号的地址才是安全的。cudartcudart.liblibcudart.acudart.dlllibcudart.socudart.dllcudart.so

它的所有入口点都以 为前缀。cuda

如异构编程中所述，CUDA 编程模型假设系统由主机和设备组成，每个设备都有自己独立的内存。设备内存概述了用于管理设备内存的运行时函数。

共享内存演示了如何使用线程层次结构中引入的共享内存来最大化性能。

页面锁定主机内存引入了页面锁定主机内存，这是将内核执行与主机和设备内存之间的数据传输重叠所必需的。

异步并发执行描述了用于在系统中的各个级别启用异步并发执行的概念和 API。

多设备系统展示了编程模型如何扩展到多个设备连接到同一主机的系统。

错误检查介绍如何正确检查运行时产生的错误。

调用堆栈提到了用于管理 CUDA C++ 调用堆栈的运行时函数。

纹理和表面内存 提供另一种访问设备内存的方法的纹理和表面内存空间;它们还公开了 GPU 纹理硬件的一个子集。

图形互操作性介绍了运行时提供的各种函数，用于与两个主要的图形 API（OpenGL 和 Direct3D）进行互操作。

3.2.1. 初始化

从 CUDA 12.0 开始，and 调用初始化运行时以及与指定设备关联的主要上下文。如果没有这些调用，运行时将隐式使用设备 0 并根据需要自行初始化以处理其他运行时 API 请求。在对运行时函数调用进行计时以及将第一次调用的错误代码解释到运行时时，需要牢记这一点。在 12.0 之前，不会初始化运行时，应用程序通常会使用无操作运行时调用来将运行时初始化与其他 API 活动隔离开来（为了计时和错误处理）。cudaInitDevice()cudaSetDevice()cudaSetDevice()cudaFree(0)

运行时为系统中的每个设备创建一个 CUDA 上下文（有关 CUDA 上下文的更多详细信息，请参阅上下文）。此上下文是此设备的主要上下文，并在第一个运行时函数时初始化，该函数需要此设备上的活动上下文。它在应用程序的所有主机线程之间共享。作为此上下文创建的一部分，如有必要，将对设备代码进行实时编译（请参阅实时编译）并加载到设备内存中。这一切都是透明的。如果需要，例如，为了实现驱动程序 API 互操作性，可以从驱动程序 API 访问设备的主要上下文，如运行时 API 和驱动程序 API 之间的互操作性中所述。

当主机线程调用时，这会破坏主机线程当前运行的设备的主要上下文（即，在设备选择中定义的当前设备）。将此设备设置为当前设备的任何主机线程进行的下一个运行时函数调用将为该设备创建新的主上下文。cudaDeviceReset()

3.2.2. 设备内存

如异构编程中所述，CUDA 编程模型假设系统由主机和设备组成，每个设备都有自己独立的内存。内核在设备内存不足的情况下运行，因此运行时提供分配、解除分配和复制设备内存以及在主机内存和设备内存之间传输数据的函数。

设备内存可以分配为线性内存或 CUDA 数组。

CUDA 数组是针对纹理获取进行了优化的不透明内存布局。它们在纹理和表面内存中进行了描述。

线性内存在单个统一的地址空间中分配，这意味着单独分配的实体可以通过指针相互引用，例如，在二叉树或链表中。地址空间的大小取决于主机系统 （CPU） 和所用 GPU 的计算能力：

 线性内存通常使用 和释放 使用进行分配，主机内存和设备内存之间的数据传输通常使用 。在 Kernels 的向量加法代码示例中，需要将向量从主机内存复制到设备内存中：cudaMalloc()cudaFree()cudaMemcpy()

#include <iostream>
#include <cuda_runtime.h>__global__ void VecAdd(float* A, float* B, float* C, int N)
{int i = blockDim.x * blockIdx.x + threadIdx.x;if (i < N)C[i] = A[i] + B[i];
}int main()
{int N = 1024; // Size of the vectorssize_t size = N * sizeof(float);// Allocate input vectors in host memoryfloat* h_A = (float*)malloc(size);float* h_B = (float*)malloc(size);float* h_C = (float*)malloc(size);// Initialize input vectorsfor (int i = 0; i < N; ++i){h_A[i] = static_cast<float>(i);h_B[i] = static_cast<float>(i * 2);}// Allocate vectors in device memoryfloat* d_A;cudaMalloc(&d_A, size);float* d_B;cudaMalloc(&d_B, size);float* d_C;cudaMalloc(&d_C, size);// Copy vectors from host memory to device memorycudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);// Invoke kernelint threadsPerBlock = 256;int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);// Copy result from device memory to host memorycudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);// Print result (optional)std::cout << "Result vector C:" << std::endl;for (int i = 0; i < N; ++i){std::cout << h_C[i] << " ";if ((i + 1) % 10 == 0) std::cout << std::endl; // Print in rows of 10}// Free device memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// Free host memoryfree(h_A);free(h_B);free(h_C);return 0;
}

输出：

Result vector C:
0 3 6 9 12 15 18 21 24 27 
30 33 36 39 42 45 48 51 54 57 
60 63 66 69 72 75 78 81 84 87 
90 93 96 99 102 105 108 111 114 117 
120 123 126 129 132 135 138 141 144 147 
150 153 156 159 162 165 168 171 174 177 
180 183 186 189 192 195 198 201 204 207 
210 213 216 219 222 225 228 231 234 237 
240 243 246 249 252 255 258 261 264 267 
270 273 276 279 282 285 288 291 294 297 
300 303 306 309 312 315 318 321 324 327 
330 333 336 339 342 345 348 351 354 357 
360 363 366 369 372 375 378 381 384 387 
390 393 396 399 402 405 408 411 414 417 
420 423 426 429 432 435 438 441 444 447 
450 453 456 459 462 465 468 471 474 477 
480 483 486 489 492 495 498 501 504 507 
510 513 516 519 522 525 528 531 534 537 
540 543 546 549 552 555 558 561 564 567 
570 573 576 579 582 585 588 591 594 597 
600 603 606 609 612 615 618 621 624 627 
630 633 636 639 642 645 648 651 654 657 
660 663 666 669 672 675 678 681 684 687 
690 693 696 699 702 705 708 711 714 717 
720 723 726 729 732 735 738 741 744 747 
750 753 756 759 762 765 768 771 774 777 
780 783 786 789 792 795 798 801 804 807 
810 813 816 819 822 825 828 831 834 837 
840 843 846 849 852 855 858 861 864 867 
870 873 876 879 882 885 888 891 894 897 
900 903 906 909 912 915 918 921 924 927 
930 933 936 939 942 945 948 951 954 957 
960 963 966 969 972 975 978 981 984 987 
990 993 996 999 1002 1005 1008 1011 1014 1017 
1020 1023 1026 1029 1032 1035 1038 1041 1044 1047 
1050 1053 1056 1059 1062 1065 1068 1071 1074 1077 
1080 1083 1086 1089 1092 1095 1098 1101 1104 1107 
1110 1113 1116 1119 1122 1125 1128 1131 1134 1137 
1140 1143 1146 1149 1152 1155 1158 1161 1164 1167 
1170 1173 1176 1179 1182 1185 1188 1191 1194 1197 
1200 1203 1206 1209 1212 1215 1218 1221 1224 1227 
1230 1233 1236 1239 1242 1245 1248 1251 1254 1257 
1260 1263 1266 1269 1272 1275 1278 1281 1284 1287 
1290 1293 1296 1299 1302 1305 1308 1311 1314 1317 
1320 1323 1326 1329 1332 1335 1338 1341 1344 1347 
1350 1353 1356 1359 1362 1365 1368 1371 1374 1377 
1380 1383 1386 1389 1392 1395 1398 1401 1404 1407 
1410 1413 1416 1419 1422 1425 1428 1431 1434 1437 
1440 1443 1446 1449 1452 1455 1458 1461 1464 1467 
1470 1473 1476 1479 1482 1485 1488 1491 1494 1497 
1500 1503 1506 1509 1512 1515 1518 1521 1524 1527 
1530 1533 1536 1539 1542 1545 1548 1551 1554 1557 
1560 1563 1566 1569 1572 1575 1578 1581 1584 1587 
1590 1593 1596 1599 1602 1605 1608 1611 1614 1617 
1620 1623 1626 1629 1632 1635 1638 1641 1644 1647 
1650 1653 1656 1659 1662 1665 1668 1671 1674 1677 
1680 1683 1686 1689 1692 1695 1698 1701 1704 1707 
1710 1713 1716 1719 1722 1725 1728 1731 1734 1737 
1740 1743 1746 1749 1752 1755 1758 1761 1764 1767 
1770 1773 1776 1779 1782 1785 1788 1791 1794 1797 
1800 1803 1806 1809 1812 1815 1818 1821 1824 1827 
1830 1833 1836 1839 1842 1845 1848 1851 1854 1857 
1860 1863 1866 1869 1872 1875 1878 1881 1884 1887 
1890 1893 1896 1899 1902 1905 1908 1911 1914 1917 
1920 1923 1926 1929 1932 1935 1938 1941 1944 1947 
1950 1953 1956 1959 1962 1965 1968 1971 1974 1977 
1980 1983 1986 1989 1992 1995 1998 2001 2004 2007 
2010 2013 2016 2019 2022 2025 2028 2031 2034 2037 
2040 2043 2046 2049 2052 2055 2058 2061 2064 2067 
2070 2073 2076 2079 2082 2085 2088 2091 2094 2097 
2100 2103 2106 2109 2112 2115 2118 2121 2124 2127 
2130 2133 2136 2139 2142 2145 2148 2151 2154 2157 
2160 2163 2166 2169 2172 2175 2178 2181 2184 2187 
2190 2193 2196 2199 2202 2205 2208 2211 2214 2217 
2220 2223 2226 2229 2232 2235 2238 2241 2244 2247 
2250 2253 2256 2259 2262 2265 2268 2271 2274 2277 
2280 2283 2286 2289 2292 2295 2298 2301 2304 2307 
2310 2313 2316 2319 2322 2325 2328 2331 2334 2337 
2340 2343 2346 2349 2352 2355 2358 2361 2364 2367 
2370 2373 2376 2379 2382 2385 2388 2391 2394 2397 
2400 2403 2406 2409 2412 2415 2418 2421 2424 2427 
2430 2433 2436 2439 2442 2445 2448 2451 2454 2457 
2460 2463 2466 2469 2472 2475 2478 2481 2484 2487 
2490 2493 2496 2499 2502 2505 2508 2511 2514 2517 
2520 2523 2526 2529 2532 2535 2538 2541 2544 2547 
2550 2553 2556 2559 2562 2565 2568 2571 2574 2577 
2580 2583 2586 2589 2592 2595 2598 2601 2604 2607 
2610 2613 2616 2619 2622 2625 2628 2631 2634 2637 
2640 2643 2646 2649 2652 2655 2658 2661 2664 2667 
2670 2673 2676 2679 2682 2685 2688 2691 2694 2697 
2700 2703 2706 2709 2712 2715 2718 2721 2724 2727 
2730 2733 2736 2739 2742 2745 2748 2751 2754 2757 
2760 2763 2766 2769 2772 2775 2778 2781 2784 2787 
2790 2793 2796 2799 2802 2805 2808 2811 2814 2817 
2820 2823 2826 2829 2832 2835 2838 2841 2844 2847 
2850 2853 2856 2859 2862 2865 2868 2871 2874 2877 
2880 2883 2886 2889 2892 2895 2898 2901 2904 2907 
2910 2913 2916 2919 2922 2925 2928 2931 2934 2937 
2940 2943 2946 2949 2952 2955 2958 2961 2964 2967 
2970 2973 2976 2979 2982 2985 2988 2991 2994 2997 
3000 3003 3006 3009 3012 3015 3018 3021 3024 3027 
3030 3033 3036 3039 3042 3045 3048 3051 3054 3057 
3060 3063 3066 3069

线性内存也可以通过 和 进行分配。建议将这些函数用于 2D 或 3D 数组的分配，因为它可确保适当填充分配以满足设备内存访问中所述的对齐要求，从而确保在访问行地址或在 2D 数组和设备内存的其他区域之间执行复制时获得最佳性能（使用 and 函数）。返回的音高（或步幅）必须用于访问数组元素。以下代码示例分配一个浮点值的 x 2D 数组，并演示如何在设备代码中循环访问数组元素：cudaMallocPitch()cudaMalloc3D()cudaMemcpy2D()cudaMemcpy3D()widthheight

// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);// Device code
__global__ void MyKernel(float* devPtr,size_t pitch, int width, int height)
{for (int r = 0; r < height; ++r) {float* row = (float*)((char*)devPtr + r * pitch);for (int c = 0; c < width; ++c) {float element = row[c];}}
}

以下代码示例分配浮点值的 x x 3D 数组，并演示如何在设备代码中循环访问数组元素：widthheightdepth

// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),height, depth);
cudaPitchedPtr devPitchedPtr;
cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,int width, int height, int depth)
{char* devPtr = devPitchedPtr.ptr;size_t pitch = devPitchedPtr.pitch;size_t slicePitch = pitch * height;for (int z = 0; z < depth; ++z) {char* slice = devPtr + z * slicePitch;for (int y = 0; y < height; ++y) {float* row = (float*)(slice + y * pitch);for (int x = 0; x < width; ++x) {float element = row[x];}}}
}

参考手册列出了用于在线性内存分配的线性内存、分配的线性内存和为全局或常量内存空间中声明的变量分配的内存之间复制内存的所有各种函数。cudaMalloc()cudaMallocPitch()cudaMalloc3D()

以下代码示例演示了通过运行时 API 访问全局变量的各种方法：

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

cudaGetSymbolAddress()用于检索指向为全局内存空间中声明的变量分配的内存的地址。分配的内存的大小是通过 获得的。cudaGetSymbolSize()

3.2.4. 共享内存 

如可变内存空间说明符中所述，共享内存是使用内存空间说明符分配的。__shared__

共享内存预计比全局内存快得多，如线程层次结构中所述，并在共享内存中进行了详细说明。它可以用作暂存器内存（或软件管理的缓存），以最大程度地减少来自 CUDA 块的全局内存访问，如以下矩阵乘法示例所示。

以下代码示例是矩阵乘法的简单实现，它不利用共享内存。每个线程读取一行 A 和一列 B，并计算 C 的相应元素，如图 8 所示。因此，A是读取B.从全局内存中读取的宽度时间，B是读取A.高度时间。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {int width;int height;float* elements;
} Matrix;// Thread block size
#define BLOCK_SIZE 16// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{// Load A and B to device memoryMatrix d_A;d_A.width = A.width; d_A.height = A.height;size_t size = A.width * A.height * sizeof(float);cudaMalloc(&d_A.elements, size);cudaMemcpy(d_A.elements, A.elements, size,cudaMemcpyHostToDevice);Matrix d_B;d_B.width = B.width; d_B.height = B.height;size = B.width * B.height * sizeof(float);cudaMalloc(&d_B.elements, size);cudaMemcpy(d_B.elements, B.elements, size,cudaMemcpyHostToDevice);// Allocate C in device memoryMatrix d_C;d_C.width = C.width; d_C.height = C.height;size = C.width * C.height * sizeof(float);cudaMalloc(&d_C.elements, size);// Invoke kerneldim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);// Read C from device memorycudaMemcpy(C.elements, d_C.elements, size,cudaMemcpyDeviceToHost);// Free device memorycudaFree(d_A.elements);cudaFree(d_B.elements);cudaFree(d_C.elements);
}// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{// Each thread computes one element of C// by accumulating results into Cvaluefloat Cvalue = 0;int row = blockIdx.y * blockDim.y + threadIdx.y;int col = blockIdx.x * blockDim.x + threadIdx.x;for (int e = 0; e < A.width; ++e)Cvalue += A.elements[row * A.width + e]* B.elements[e * B.width + col];C.elements[row * C.width + col] = Cvalue;
}

以下代码示例是矩阵乘法的实现，它利用了共享内存。在此实现中，每个线程块负责计算 C 的一个方形子矩阵 Csub，块中的每个线程负责计算 Csub 的一个元素。如图 9 所示，Csub 等于两个矩形矩阵的乘积：维度 （A.width， block_size） 的 A 子矩阵与 Csub 具有相同的行索引，维度 （block_size， A.width） 的 B 子矩阵与 Csub 具有相同的列索引.为了适应设备的资源，这两个矩形矩阵被划分为尽可能多的维数为 block_size 的方阵矩阵，并将 Csub 计算为这些方阵矩阵的乘积之和。这些产品中的每一个都是通过首先将两个相应的方阵从全局内存加载到共享内存中来执行的，其中一个线程加载每个矩阵的一个元素，然后让每个线程计算乘积的一个元素。每个线程将这些产品中的每一个的结果累积到一个寄存器中，一旦完成，将结果写入全局内存。

通过以这种方式阻止计算，我们利用了快速共享内存并节省了大量的全局内存带宽，因为 A 只是从全局内存中读取 （B.width / block_size） 次，而 B 是读取 （A.height / block_size） 次。

上一个代码示例中的 Matrix 类型使用步幅字段进行了增强，以便可以使用相同的类型有效地表示子矩阵。__device__函数用于获取和设置元素，并从矩阵构建任何子矩阵。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {int width;int height;int stride;float* elements;
} Matrix;
// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{return A.elements[row * A.stride + col];
}
// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col,float value)
{A.elements[row * A.stride + col] = value;
}
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A__device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{Matrix Asub;Asub.width    = BLOCK_SIZE;Asub.height   = BLOCK_SIZE;Asub.stride   = A.stride;Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row+ BLOCK_SIZE * col];return Asub;
}
// Thread block size
#define BLOCK_SIZE 16
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{// Load A and B to device memoryMatrix d_A;d_A.width = d_A.stride = A.width; d_A.height = A.height;size_t size = A.width * A.height * sizeof(float);cudaMalloc(&d_A.elements, size);cudaMemcpy(d_A.elements, A.elements, size,cudaMemcpyHostToDevice);Matrix d_B;d_B.width = d_B.stride = B.width; d_B.height = B.height;size = B.width * B.height * sizeof(float);cudaMalloc(&d_B.elements, size);cudaMemcpy(d_B.elements, B.elements, size,cudaMemcpyHostToDevice);// Allocate C in device memoryMatrix d_C;d_C.width = d_C.stride = C.width; d_C.height = C.height;size = C.width * C.height * sizeof(float);cudaMalloc(&d_C.elements, size);// Invoke kerneldim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);// Read C from device memorycudaMemcpy(C.elements, d_C.elements, size,cudaMemcpyDeviceToHost);// Free device memorycudaFree(d_A.elements);cudaFree(d_B.elements);cudaFree(d_C.elements);
}
// Matrix multiplication kernel called by MatMul()__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{// Block row and columnint blockRow = blockIdx.y;int blockCol = blockIdx.x;// Each thread block computes one sub-matrix Csub of CMatrix Csub = GetSubMatrix(C, blockRow, blockCol);// Each thread computes one element of Csub// by accumulating results into Cvaluefloat Cvalue = 0;// Thread row and column within Csubint row = threadIdx.y;int col = threadIdx.x;// Loop over all the sub-matrices of A and B that are// required to compute Csub// Multiply each pair of sub-matrices together// and accumulate the resultsfor (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {// Get sub-matrix Asub of AMatrix Asub = GetSubMatrix(A, blockRow, m);// Get sub-matrix Bsub of BMatrix Bsub = GetSubMatrix(B, m, blockCol);// Shared memory used to store Asub and Bsub respectively__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];// Load Asub and Bsub from device memory to shared memory// Each thread loads one element of each sub-matrixAs[row][col] = GetElement(Asub, row, col);Bs[row][col] = GetElement(Bsub, row, col);// Synchronize to make sure the sub-matrices are loaded// before starting the computation__syncthreads();// Multiply Asub and Bsub togetherfor (int e = 0; e < BLOCK_SIZE; ++e)Cvalue += As[row][e] * Bs[e][col];// Synchronize to make sure that the preceding// computation is done before loading two new// sub-matrices of A and B in the next iteration__syncthreads();}// Write Csub to device memory// Each thread writes one elementSetElement(Csub, row, col, Cvalue);
}

3.2.5. 分布式共享内存

计算能力 9.0 中引入的线程块集群为线程块集群中的线程提供了访问集群中所有参与线程块的共享内存的能力。这种分区的共享内存称为分布式共享内存，对应的地址空间称为分布式共享内存地址空间。属于线程块集群的线程，可以在分布式地址空间中读取、写入或执行原子操作，无论该地址是属于本地线程块还是远程线程块。无论内核是否使用分布式共享内存，共享内存大小规格，静态或动态仍然是每个线程块。分布式共享内存的大小就是每个集群的线程块数乘以每个线程块的共享内存大小。

访问分布式共享内存中的数据需要所有线程块都存在。用户可以保证所有线程块都已开始使用集群组 API 执行。 用户还需要确保所有分布式共享内存操作都发生在线程块退出之前，例如，如果远程线程块试图读取给定线程块的共享内存，用户需要确保远程线程块读取的共享内存在退出之前已经完成。cluster.sync()

CUDA 提供了一种访问分布式共享内存的机制，应用程序可以从利用其功能中受益。让我们看一下一个简单的直方图计算，以及如何使用线程块集群在 GPU 上优化它。计算直方图的标准方法是在每个线程块的共享内存中进行计算，然后执行全局内存原子分析。这种方法的一个限制是共享内存容量。一旦直方图条柱不再适合共享内存，用户就需要直接计算直方图，从而计算全局内存中的原子。对于分布式共享内存，CUDA提供了一个中间步骤，其中根据直方图箱的大小，直方图可以直接在共享内存、分布式共享内存或全局内存中计算。

下面的 CUDA 内核示例展示了如何计算共享内存或分布式共享内存中的直方图，具体取决于直方图条柱的数量。

#include <cooperative_groups.h>// Distributed Shared memory histogram kernel
__global__ void clusterHist_kernel(int *bins, const int nbins, const int bins_per_block, const int *__restrict__ input,size_t array_size)
{extern __shared__ int smem[];namespace cg = cooperative_groups;int tid = cg::this_grid().thread_rank();// Cluster initialization, size and calculating local bin offsets.cg::cluster_group cluster = cg::this_cluster();unsigned int clusterBlockRank = cluster.block_rank();int cluster_size = cluster.dim_blocks().x;for (int i = threadIdx.x; i < bins_per_block; i += blockDim.x){smem[i] = 0; //Initialize shared memory histogram to zeros}// cluster synchronization ensures that shared memory is initialized to zero in// all thread blocks in the cluster. It also ensures that all thread blocks// have started executing and they exist concurrently.cluster.sync();for (int i = tid; i < array_size; i += blockDim.x * gridDim.x){int ldata = input[i];//Find the right histogram bin.int binid = ldata;if (ldata < 0)binid = 0;else if (ldata >= nbins)binid = nbins - 1;//Find destination block rank and offset for computing//distributed shared memory histogramint dst_block_rank = (int)(binid / bins_per_block);int dst_offset = binid % bins_per_block;//Pointer to target block shared memoryint *dst_smem = cluster.map_shared_rank(smem, dst_block_rank);//Perform atomic update of the histogram binatomicAdd(dst_smem + dst_offset, 1);}// cluster synchronization is required to ensure all distributed shared// memory operations are completed and no thread block exits while// other thread blocks are still accessing distributed shared memorycluster.sync();// Perform global memory histogram, using the local distributed memory histogramint *lbins = bins + cluster.block_rank() * bins_per_block;for (int i = threadIdx.x; i < bins_per_block; i += blockDim.x){atomicAdd(&lbins[i], smem[i]);}
}

 上述内核可以在运行时启动，集群大小取决于所需的分布式共享内存量。如果直方图足够小，可以只容纳一个块的共享内存，用户可以启动集群大小为 1 的内核。下面的代码片段展示了如何根据共享内存要求动态启动集群内核。

// Launch via extensible launch
{cudaLaunchConfig_t config = {0};config.gridDim = array_size / threads_per_block;config.blockDim = threads_per_block;// cluster_size depends on the histogram size.// ( cluster_size == 1 ) implies no distributed shared memory, just thread block local shared memoryint cluster_size = 2; // size 2 is an example hereint nbins_per_block = nbins / cluster_size;//dynamic shared memory size is per block.//Distributed shared memory size =  cluster_size * nbins_per_block * sizeof(int)config.dynamicSmemBytes = nbins_per_block * sizeof(int);CUDA_CHECK(::cudaFuncSetAttribute((void *)clusterHist_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, config.dynamicSmemBytes));cudaLaunchAttribute attribute[1];attribute[0].id = cudaLaunchAttributeClusterDimension;attribute[0].val.clusterDim.x = cluster_size;attribute[0].val.clusterDim.y = 1;attribute[0].val.clusterDim.z = 1;config.numAttrs = 1;config.attrs = attribute;cudaLaunchKernelEx(&config, clusterHist_kernel, bins, nbins, nbins_per_block, input, array_size);
}

3.2.6. 页面锁定主机内存

运行时提供的功能允许使用页面锁定（也称为固定）主机内存（而不是由分配的常规可分页主机内存）：malloc()

• cudaHostAlloc()分配和释放页面锁定的主机内存;cudaFreeHost()

• cudaHostRegister()page-locks分配的内存范围（有关限制，请参阅参考手册）。malloc()

使用页面锁定主机内存有几个好处：

• 对于某些设备，页面锁定主机内存和设备内存之间的复制可以与内核执行同时执行，如异步并发执行中所述。

• 在某些设备上，页面锁定的主机内存可以映射到设备的地址空间，从而无需将其复制到设备内存或从设备内存中复制，如映射内存中所述。

• 在具有前端总线的系统上，如果主机内存被分配为页面锁定，则主机内存和设备内存之间的带宽会更高，如果主机内存被分配为写入组合，则带宽会更高，如写入组合内存中所述。

简单的零拷贝 CUDA 示例附带了有关页面锁定内存 API 的详细文档。

3.2.6.1. 便携式内存 

 页面锁定内存块可以与系统中的任何设备结合使用（有关多设备系统的更多详细信息，请参阅多设备系统），但默认情况下，使用上述页面锁定内存的好处仅在与分配块时处于当前状态的设备一起使用（并且所有设备共享相同的统一地址空间， 如果有，如统一虚拟地址空间中所述）。为了使这些优势适用于所有设备，需要通过将标志传递给来分配块，或者通过将标志传递给 来锁定页面。cudaHostAllocPortablecudaHostAlloc()cudaHostRegisterPortablecudaHostRegister()

3.2.6.2. 写合并内存 

默认情况下，页面锁定的主机内存被分配为可缓存。可以选择性地将其分配为写入组合，而是通过将标志传递给 。写入组合内存可释放主机的 L1 和 L2 缓存资源，从而为应用程序的其余部分提供更多缓存。此外，在通过 PCI Express 总线进行传输期间，写入组合内存不会被窥探，这可以将传输性能提高多达 40%。cudaHostAllocWriteCombinedcudaHostAlloc()

从主机的写入组合内存中读取速度非常慢，因此写入合并内存通常应用于主机仅写入的内存。

应避免在 WC 内存上使用 CPU 原子指令，因为并非所有 CPU 实现都保证该功能。

3.2.6.3. 映射内存

也可以通过将标志传递给 或将标志传递给 来将页面锁定的主机内存块映射到设备的地址空间。因此，这样的块通常有两个地址：一个位于主机内存中，由 or 返回，另一个位于设备内存中，可以使用该地址进行检索，然后用于从内核内部访问该块。唯一的例外是，当主机和设备使用统一地址空间时，分配了指针，如统一虚拟地址空间中所述。cudaHostAllocMappedcudaHostAlloc()cudaHostRegisterMappedcudaHostRegister()cudaHostAlloc()malloc()cudaHostGetDevicePointer()cudaHostAlloc()

直接从内核内部访问主机内存并不能提供与设备内存相同的带宽，但确实有一些优点：

• 无需在设备内存中分配一个块，并在此块和主机内存中的块之间复制数据;数据传输是根据内核的需要隐式执行的;

• 无需使用流（请参阅并发数据传输）来将数据传输与内核执行重叠;内核发起的数据传输会自动与内核执行重叠。

但是，由于映射的页面锁定内存在主机和设备之间共享，因此应用程序必须使用流或事件同步内存访问（请参阅异步并发执行），以避免任何潜在的先写后读、先读后写或先写后写的危险。

为了能够检索指向任何映射的页面锁定内存的设备指针，必须在执行任何其他 CUDA 调用之前通过使用标志调用来启用页面锁定内存映射。否则，将返回错误。cudaSetDeviceFlags()cudaDeviceMapHostcudaHostGetDevicePointer()

cudaHostGetDevicePointer()如果设备不支持映射的页面锁定主机内存，也会返回错误。应用程序可以通过检查设备属性（请参阅设备枚举）来查询此功能，对于支持映射页面锁定主机内存的设备，该属性等于 1。canMapHostMemory

请注意，从主机或其他设备的角度来看，在映射的页面锁定内存上运行的原子函数（请参阅原子函数）不是原子函数。

另请注意，CUDA 运行时要求从主机和其他设备的角度来看，将从设备启动的 1 字节、2 字节、4 字节和 8 字节自然对齐的负载和存储保留为单一访问。在某些平台上，原子到内存可能会被硬件分解为单独的加载和存储操作。这些组件加载和存储操作对保留自然对齐的访问具有相同的要求。例如，CUDA 运行时不支持 PCI Express 总线拓扑，其中 PCI Express 桥接器将 8 字节自然对齐的写入分成两个 4 字节的写入，在设备和主机之间。

3.2.7. 内存同步域

3.2.7.1. 内存栅栏干扰

某些 CUDA 应用程序可能会看到性能下降，因为内存围栏/刷新操作等待的事务数量超过了 CUDA 内存一致性模型所需的事务数量。

__managed__ int x = 0;
__device__  cuda::atomic<int, cuda::thread_scope_device> a(0);
__managed__ cuda::atomic<int, cuda::thread_scope_system> b(0);

x = 1;
a = 1;

while (a != 1) ;
assert(x == 1);
b = 1;

while (b != 1) ;
assert(x == 1);

请看上面的例子。CUDA 内存一致性模型保证断言的条件为 true，因此在从线程 2 写入之前，线程 1 的写入必须对线程 3 可见。xb

释放和获取提供的内存排序仅足以使线程 2 可见，而不是线程 3，因为它是设备范围的操作。因此，release 和 acquire 提供的系统范围排序需要确保不仅从线程 2 本身发出的写入对线程 3 可见，而且从线程 2 可见的其他线程的写入也可见。这被称为累积性。由于 GPU 在执行时无法知道哪些写入在源级别被保证是可见的，哪些写入只是通过偶然的时间可见，因此它必须为飞行中的内存操作撒下一张保守的广网。axb

这有时会导致干扰：由于 GPU 正在等待内存操作，因此在源级别不需要这样做，因此隔离/刷新可能需要更长的时间。

请注意，围栏可能在代码中显式地作为内部函数或原子出现，如示例中所示，或者隐式地在任务边界处实现同步关系。

一个常见的例子是，当一个内核在本地 GPU 内存中执行计算，而一个并行内核（例如来自 NCCL 的内核）正在与对等体执行通信。完成后，本地内核将隐式刷新其写入操作，以满足与下游工作的任何同步关系。这可能会不必要地全部或部分等待来自通信内核的较慢的 nvlink 或 PCIe 写入。

3.2.7.2. 将流量与域隔离

从 Hopper 架构 GPU 和 CUDA 12.0 开始，内存同步域功能提供了一种减轻此类干扰的方法。作为代码的明确帮助的交换，GPU 可以减少围栏操作造成的网络投射。每次内核启动都会被赋予一个域 ID。写入和围栏都用 ID 标记，而围栏只会对匹配围栏域的写入进行排序。在并发计算与通信示例中，通信内核可以放置在不同的域中。

使用域时，代码必须遵守以下规则，即在同一 GPU 上的不同域之间排序或同步需要系统范围隔离。在域中，设备范围的隔离仍然足够了。这对于累积性是必要的，因为一个内核的写入不会被另一个域中的内核发出的栅栏所包含。从本质上讲，通过确保提前将跨域流量刷新到系统范围来满足累积性。

请注意，这将修改 的定义。但是，由于内核将默认为域 0（如下所述），因此可以保持向后兼容性。thread_scope_device

3.2.7.3. 在 CUDA 中使用域

可以通过新的启动属性和 来访问域。前者在逻辑域和 之间进行选择，后者提供从逻辑域到物理域的映射。远程域用于执行远程内存访问的内核，以便将其内存流量与本地内核隔离开来。但是请注意，特定域的选择不会影响内核可以合法执行的内存访问。cudaLaunchAttributeMemSyncDomaincudaLaunchAttributeMemSyncDomainMapcudaLaunchMemSyncDomainDefaultcudaLaunchMemSyncDomainRemote

可以通过 device 属性查询域计数。Hopper 有 4 个域。为了便于移植代码，域功能可以在所有设备上使用，CUDA 将在 Hopper 之前报告计数为 1。cudaDevAttrMemSyncDomainCount

拥有逻辑域可以简化应用程序组合。在堆栈的较低级别启动单个内核（例如从 NCCL 启动）可以选择语义逻辑域，而无需关注周围的应用程序架构。更高级别的可以使用映射来引导逻辑域。如果未设置逻辑域的默认值，则该值为默认域，默认映射是将默认域映射到 0，将远程域映射到 1（在具有 1 个以上域的 GPU 上）。特定库可能会在 CUDA 12.0 及更高版本中使用远程域标记启动;例如，NCCL 2.16 将这样做。总之，这为开箱即用的常见应用程序提供了一种有益的使用模式，无需在其他组件、框架或应用程序级别更改代码。另一种使用模式，例如在使用 nvshmem 的应用程序中或没有明确分离内核类型的应用程序中，可能是对并行流进行分区。流 A 可以将两个逻辑域映射到物理域 0，将流 B 映射到 1，依此类推。

// Example of launching a kernel with the remote logical domain
cudaLaunchAttribute domainAttr;
domainAttr.id = cudaLaunchAttrMemSyncDomain;
domainAttr.val = cudaLaunchMemSyncDomainRemote;
cudaLaunchConfig_t config;
// Fill out other config fields
config.attrs = &domainAttr;
config.numAttrs = 1;
cudaLaunchKernelEx(&config, myKernel, kernelArg1, kernelArg2...);

// Example of setting a mapping for a stream
// (This mapping is the default for streams starting on Hopper if not
// explicitly set, and provided for illustration)
cudaLaunchAttributeValue mapAttr;
mapAttr.memSyncDomainMap.default_ = 0;
mapAttr.memSyncDomainMap.remote = 1;
cudaStreamSetAttribute(stream, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);

// Example of mapping different streams to different physical domains, ignoring
// logical domain settings
cudaLaunchAttributeValue mapAttr;
mapAttr.memSyncDomainMap.default_ = 0;
mapAttr.memSyncDomainMap.remote = 0;
cudaStreamSetAttribute(streamA, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);
mapAttr.memSyncDomainMap.default_ = 1;
mapAttr.memSyncDomainMap.remote = 1;
cudaStreamSetAttribute(streamB, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);

与其他启动属性一样，这些属性在 CUDA 流、单个启动 using 和 CUDA 图中的内核节点上统一公开。如上所述，典型的用途是在流级别设置映射，在启动级别设置逻辑域（或将流使用的一部分括起来）。cudaLaunchKernelEx

在流捕获期间，这两个属性都会复制到图形节点。图形从节点本身获取这两个属性，本质上是一种指定物理域的间接方式。在启动图形的流上设置的域相关属性不会在图形的执行中使用。

3.2.8. 异步并发执行 

CUDA 将以下操作公开为可以并发运行的独立任务：

• 在主机上计算;

• 在设备上进行计算;

• 内存从主机传输到设备;

• 内存从设备传输到主机;

• 在给定设备的内存内进行内存传输;

• 设备之间的内存传输。

这些操作之间实现的并发级别将取决于设备的功能集和计算能力，如下所述。

 3.2.8.1. 主机和设备之间的并发执行