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 memory
    float *h_A = new float[N];
    float *h_B = new float[N];
    float *h_C = new float[N];

    // Initialize host vectors
    for (int i = 0; i < N; ++i)
    {
        h_A[i] = i * 1.0f; // Example values
        h_B[i] = i * 2.0f;
    }

    // Allocate device memory
    float *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 device
    cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);

    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(d_A, d_B, d_C);

    // Copy result from device to host
    cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

    // Print result
    std::cout << "Result vector C:" << std::endl;
    for (int i = 0; i < N; ++i)
    {
        std::cout << h_C[i] << " ";
    }
    std::cout << std::endl;

    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    delete[] 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 access
        C[index] = A[index] + B[index];
    }
}

int main()
{
    // Allocate host memory
    float *h_A = new float[N * N];
    float *h_B = new float[N * N];
    float *h_C = new float[N * N];

    // Initialize host matrices
    for (int i = 0; i < N; ++i)
    {
        for (int j = 0; j < N; ++j)
        {
            h_A[i * N + j] = static_cast<float>(i + j); // Example initialization
            h_B[i * N + j] = static_cast<float>(i - j); // Example initialization
        }
    }

    // Allocate device memory
    float *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 device
    cudaMemcpy(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 threads
    dim3 threadsPerBlock(N, N);
    MatAdd<<<1, threadsPerBlock>>>(d_A, d_B, d_C, N);

    // Copy result from device to host
    cudaMemcpy(h_C, d_C, N * N * sizeof(float), cudaMemcpyDeviceToHost);

    // Print result
    std::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 memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    delete[] 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 access
        C[index] = A[index] + B[index];
    }
}

int main()
{
    // Allocate host memory
    float *h_A = new float[N * N];
    float *h_B = new float[N * N];
    float *h_C = new float[N * N];

    // Initialize host matrices
    for (int i = 0; i < N; ++i)
    {
        for (int j = 0; j < N; ++j)
        {
            h_A[i * N + j] = static_cast<float>(i + j); // Example initialization
            h_B[i * N + j] = static_cast<float>(i - j); // Example initialization
        }
    }

    // Allocate device memory
    float *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 device
    cudaMemcpy(d_A, h_A, N * N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, N * N * sizeof(float), cudaMemcpyHostToDevice);

    // Kernel invocation
    dim3 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 host
    cudaMemcpy(h_C, d_C, N * N * sizeof(float), cudaMemcpyDeviceToHost);

    // Print result
    std::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 memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    delete[] 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::thread_scope::thread_scope_thread`	只有启动异步操作的 CUDA 线程才会同步。
`cuda::thread_scope::thread_scope_block`	与启动线程同步的同一线程块中的所有或任何 CUDA 线程。
`cuda::thread_scope::thread_scope_device`	与启动线程相同的 GPU 设备中的所有或任何 CUDA 线程都会同步。
`cuda::thread_scope::thread_scope_system`	与启动线程相同的系统中的所有或任何 CUDA 或 CPU 线程同步。

这些线程作用域是作为 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 Runtime

runtime是在库中实现的，该库通过 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 的计算能力：

表 1 线性内存地址空间
	x86_64 （AMD64）	电源（ppc64le）	ARM64系列
最高计算能力 5.3 （Maxwell）	40位	40位	40位
计算能力 6.0 （Pascal）或更高版本	高达 47 位	高达 49 位	高达 48 位

线性内存通常使用和释放使用进行分配，主机内存和设备内存之间的数据传输通常使用。在 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 vectors
    size_t size = N * sizeof(float);

    // Allocate input vectors in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    float* h_C = (float*)malloc(size);

    // Initialize input vectors
    for (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 memory
    float* 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 memory
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // Invoke kernel
    int 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 memory
    cudaMemcpy(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 memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    // Free host memory
    free(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 memory
    Matrix 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 memory
    Matrix 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 kernel
    dim3 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 memory
    cudaMemcpy(C.elements, d_C.elements, size,
               cudaMemcpyDeviceToHost);

    // Free device memory
    cudaFree(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 Cvalue
    float 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 memory
    Matrix 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 memory
    Matrix 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 kernel
    dim3 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 memory
    cudaMemcpy(C.elements, d_C.elements, size,
               cudaMemcpyDeviceToHost);
    // Free device memory
    cudaFree(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 column
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;
    // Each thread block computes one sub-matrix Csub of C
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
    // Each thread computes one element of Csub
    // by accumulating results into Cvalue
    float Cvalue = 0;
    // Thread row and column within Csub
    int 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 results
    for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {
        // Get sub-matrix Asub of A
        Matrix Asub = GetSubMatrix(A, blockRow, m);
        // Get sub-matrix Bsub of B
        Matrix 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-matrix
        As[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 together
        for (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 element
    SetElement(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 histogram
    int dst_block_rank = (int)(binid / bins_per_block);
    int dst_offset = binid % bins_per_block;

    //Pointer to target block shared memory
    int *dst_smem = cluster.map_shared_rank(smem, dst_block_rank);

    //Perform atomic update of the histogram bin
    atomicAdd(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 memory
  cluster.sync();

  // Perform global memory histogram, using the local distributed memory histogram
  int *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 memory
  int cluster_size = 2; // size 2 is an example here
  int 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()

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

对于某些设备，页面锁定主机内存和设备内存之间的复制可以与内核执行同时执行，如异步并发执行中所述。
在某些设备上，页面锁定的主机内存可以映射到设备的地址空间，从而无需将其复制到设备内存或从设备内存中复制，如映射内存中所述。
在具有前端总线的系统上，如果主机内存被分配为页面锁定，则主机内存和设备内存之间的带宽会更高，如果主机内存被分配为写入组合，则带宽会更高，如写入组合内存中所述。

页面锁定的主机内存不会缓存在非 I/O 相干的 Tegra 设备上。此外，在非 I/O 相干 Tegra 设备上不受支持。cudaHostRegister()

简单的零拷贝 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);

线程 1 （SM）

x = 1;
a = 1;

线程 2 （SM）

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

线程 3 （CPU）

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. 主机和设备之间的并发执行

通过异步库函数促进并发主机执行，这些函数在设备完成请求的任务之前将控制权返回给主机线程。使用异步调用，当适当的设备资源可用时，可以将许多设备操作排成一列，以便由 CUDA 驱动程序执行。这减轻了主机线程管理设备的大部分责任，使其可以自由地执行其他任务。以下设备操作相对于主机是异步的：

内核启动;
内存在单个设备的内存中复制;
将 64 KB 或更小的内存块从主机复制到设备;
由后缀为Async;
内存设置函数调用。

程序员可以通过将环境变量设置为 1 来全局禁用系统上运行的所有 CUDA 应用程序的内核启动异步性。此功能仅用于调试目的，不应用作使生产软件可靠运行的方法。CUDA_LAUNCH_BLOCKING

如果通过分析器（Nsight、Visual Profiler）收集硬件计数器，则内核启动是同步的，除非启用了并发内核分析。如果内存副本涉及未锁定页面的主机内存，则它们也可能是同步的。Async

3.2.8.2. 并发内核执行

一些计算能力为 2.x 及更高版本的设备可以同时执行多个内核。应用程序可以通过检查设备属性（请参阅设备枚举）来查询此功能，对于支持此功能的设备，该属性等于 1。concurrentKernels

设备可以并发执行的最大内核启动次数取决于其计算能力，如表 21 所示。

一个 CUDA 上下文中的内核不能与另一个 CUDA 上下文中的内核同时执行。GPU 可能会进行时间切片，以提供对每个上下文的转发进度。如果用户想在 SM 上同时运行多个进程的内核，则必须启用 MPS。

使用许多纹理或大量本地内存的内核不太可能与其他内核同时执行。

3.2.8.3. 数据传输和内核执行的重叠

某些设备可以在执行内核的同时执行到 GPU 或从 GPU 执行异步内存复制。应用程序可以通过检查设备属性（请参阅设备枚举）来查询此功能，对于支持此功能的设备，该属性大于零。如果副本中涉及主机内存，则必须将其页面锁定。asyncEngineCount

还可以在执行内核（在支持设备属性的设备上）和/或与设备之间的副本（对于支持该属性的设备）同时执行设备内复制。设备内复制是使用标准内存复制功能启动的，目标地址和源地址位于同一设备上。concurrentKernelsasyncEngineCount

3.2.8.4. 并发数据传输

某些计算能力为 2.x 及更高版本的设备可能会与设备之间的副本重叠。应用程序可以通过检查设备属性（请参阅设备枚举）来查询此功能，对于支持此功能的设备，该属性等于 2。为了重叠，传输中涉及的任何主机内存都必须进行页面锁定。asyncEngineCount

3.2.8.5. 流

应用程序通过流管理上述并发操作。流是按顺序执行的一系列命令（可能由不同的主机线程发出）。另一方面，不同的流可能会彼此之间或同时无序地执行其命令;此行为无法得到保证，因此不应依赖其正确性（例如，内核间通信未定义）。当满足命令的所有依赖项时，在流上发出的命令可能会执行。依赖项可以是以前在同一流上启动的命令，也可以是来自其他流的依赖项。成功完成同步调用可以保证启动的所有命令都已完成。

3.2.8.5.1. 流的创建和销毁

流是通过创建流对象并将其指定为一系列内核启动和主机设备内存副本的流参数来定义的。以下代码示例创建两个流，并分配一个 in page-locked memory 的数组。<->hostPtrfloat

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
    cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

以下代码示例将这些流中的每一个定义为从主机到设备的一个内存副本、一个内核启动和一个从设备到主机的内存副本的序列：

for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel <<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);
}

每个流将其输入数组的一部分复制到设备内存中的数组，通过调用在设备上进行处理，并将结果复制回的同一部分。重叠行为描述了此示例中流如何重叠，具体取决于设备的功能。请注意，必须指向页面锁定的主机内存，才会发生任何重叠。hostPtrinputDevPtrinputDevPtrMyKernel()outputDevPtrhostPtrhostPtr

流是通过调用来释放的。cudaStreamDestroy()

for (int i = 0; i < 2; ++i)
    cudaStreamDestroy(stream[i]);

如果设备在被调用时仍在流中执行工作，则该函数将立即返回，并且在设备完成流中的所有工作后，与流关联的资源将自动释放。cudaStreamDestroy()

3.2.8.5.2. 默认流

如果内核启动和主机设备内存副本未指定任何流参数，或者等效地将流参数设置为零，则将向默认流发出。因此，它们是按顺序执行的。<->

对于使用编译标志编译的代码（或在包含 CUDA 标头（和）之前定义巨集）的代码，默认流是常规流，每个主机线程都有自己的默认流。--default-stream per-threadCUDA_API_PER_THREAD_DEFAULT_STREAMcuda.hcuda_runtime.h

#define CUDA_API_PER_THREAD_DEFAULT_STREAM 1当代码由 AS 隐式包含在翻译单元的顶部编译时，不能用于启用此行为。在这种情况下，需要使用编译标志，或者需要使用编译器标志定义巨集。nvccnvcccuda_runtime.h--default-stream per-threadCUDA_API_PER_THREAD_DEFAULT_STREAM-DCUDA_API_PER_THREAD_DEFAULT_STREAM=1

对于使用编译标志编译的代码，默认流是一个称为 NULL 流的特殊流，每个设备都有一个用于所有主机线程的 NULL 流。NULL 流很特殊，因为它会导致隐式同步，如隐式同步中所述。--default-stream legacy

对于在未指定编译标志的情况下编译的代码，假定为默认值。--default-stream--default-stream legacy

。。。。。省略

3.2.8.8. 事件

runtime还提供了一种方法来密切监视设备的进度，并通过让应用程序在程序中的任何点异步记录事件，并在这些事件完成时进行查询，从而执行准确的计时。当事件之前的所有任务（或可选地，给定流中的所有命令）都已完成时，事件已完成。流 0 中的事件在完成所有流中的所有先前任务和命令后完成。

3.2.8.8.1. 事件的创建和销毁

以下代码示例创建两个事件：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

它们以这种方式被销毁：

cudaEventDestroy(start);
cudaEventDestroy(stop);

3.2.8.8.2. 经过的时间

在 Creation and Destruction 中创建的事件可用于按以下方式对 Creation 和 Destruction 的代码示例进行计时：

cudaEventRecord(start, 0);
for (int i = 0; i < 2; ++i) {
    cudaMemcpyAsync(inputDev + i * size, inputHost + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
    MyKernel<<<100, 512, 0, stream[i]>>>
               (outputDev + i * size, inputDev + i * size, size);
    cudaMemcpyAsync(outputHost + i * size, outputDev + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);
}
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);

3.2.8.9. 同步调用

调用同步函数时，在设备完成请求的任务之前，不会将控制权返回给主机线程。在主机线程执行任何其他 CUDA 调用之前，可以通过使用一些特定标志进行调用来指定主机线程是否会产生、阻塞或旋转（有关详细信息，请参阅参考手册）。cudaSetDeviceFlags()

3.2.10. 统一虚拟地址空间

当应用程序作为 64 位进程运行时，主机和计算能力为 2.0 和更高的所有设备都使用单个地址空间。通过 CUDA API 调用进行的所有主机内存分配以及受支持设备上的所有设备内存分配都在此虚拟地址范围内。因此：

在通过 CUDA 分配的主机上，或者在使用统一地址空间的任何设备上，任何内存的位置都可以通过使用指针的值来确定。cudaPointerGetAttributes()
当复制到使用统一地址空间的任何设备的内存或从内存中复制时，可以设置的参数 to 以确定指针的位置。这也适用于未通过 CUDA 分配的主机指针，只要当前设备使用统一寻址即可。cudaMemcpyKindcudaMemcpy*()cudaMemcpyDefault
分配 via 可以在使用统一地址空间的所有设备上自动移植（请参阅可移植内存），并且 by 返回的指针可以直接从这些设备上运行的内核中使用（即，无需获取设备指针 via，如映射内存中所述。cudaHostAlloc()cudaHostAlloc()cudaHostGetDevicePointer()

应用程序可以通过检查设备属性（请参阅设备枚举）是否等于 1 来查询是否为特定设备使用了统一地址空间。unifiedAddressing

3.2.11. 进程间通信

由主机线程创建的任何设备内存指针或事件句柄都可以被同一进程中的任何其他线程直接引用。但是，它在此进程之外是无效的，因此不能由属于不同进程的线程直接引用。

若要在进程之间共享设备内存指针和事件，应用程序必须使用进程间通信 API，参考手册中对此进行了详细说明。IPC API 仅支持 Linux 上的 64 位进程以及计算能力为 2.0 及更高版本的设备。请注意，分配不支持使用 IPC API。cudaMallocManaged

使用此 API，应用程序可以使用 use 获取给定设备内存指针的 IPC 句柄，将其传递给使用标准 IPC 机制的另一个进程（例如，进程间共享内存或文件），并用于从 IPC 句柄中检索设备指针，该指针是此其他进程中的有效指针。可以使用类似的入口点共享事件句柄。cudaIpcGetMemHandle()cudaIpcOpenMemHandle()

请注意，出于性能原因，所做的分配可能会从更大的内存块中进行子分配。在这种情况下，CUDA IPC API 将共享整个底层内存块，这可能会导致其他子分配被共享，这可能会导致进程之间的信息泄露。为防止此行为，建议仅共享大小为 2MiB 对齐大小的分配。cudaMalloc()

使用 IPC API 的一个示例是，单个主进程生成一批输入数据，使数据可供多个辅助进程使用，而无需重新生成或复制。

使用 CUDA IPC 相互通信的应用程序应使用相同的 CUDA 驱动程序和运行时进行编译、链接和运行。

3.2.12. 错误检查

所有运行时函数都返回错误代码，但对于异步函数（请参阅异步并发执行），此错误代码不可能报告设备上可能发生的任何异步错误，因为函数在设备完成任务之前返回;错误代码仅报告在执行任务之前在主机上发生的错误，通常与参数验证有关;如果发生异步错误，它将被一些后续不相关的运行时函数调用报告。

因此，在某个异步函数调用之后立即检查异步错误的唯一方法是在调用后立即进行同步，方法是调用（或使用异步并发执行中所述的任何其他同步机制）并检查返回的错误代码。cudaDeviceSynchronize()cudaDeviceSynchronize()

运行时为每个主机线程维护一个错误变量，每次发生错误（无论是参数验证错误还是异步错误）时，该线程都会初始化为错误代码并被错误代码覆盖。返回此变量。返回此变量并将其重置为。cudaSuccesscudaPeekAtLastError()cudaGetLastError()cudaSuccess

内核启动不会返回任何错误代码，因此必须在内核启动后立即调用，以检索任何启动前的错误。为了确保在内核启动之前由调用返回的任何错误或不是由调用引起的，必须确保将运行时错误变量设置为在内核启动之前，例如，在内核启动之前调用。内核启动是异步的，因此要检查异步错误，应用程序必须在内核启动和调用 or 之间同步。cudaPeekAtLastError()cudaGetLastError()cudaPeekAtLastError()cudaGetLastError()cudaSuccesscudaGetLastError()cudaPeekAtLastError()cudaGetLastError()

请注意，该错误可能由和返回，因此不会由或报告。cudaErrorNotReadycudaStreamQuery()cudaEventQuery()cudaPeekAtLastError()cudaGetLastError()

3.2.13. 调用堆栈

在计算能力为 2.x 及更高版本的设备上，可以使用查询和设置调用堆栈的大小。cudaDeviceGetLimit()cudaDeviceSetLimit()

当调用堆栈溢出时，如果应用程序通过 CUDA 调试器（CUDA-GDB、Nsight）运行，则内核调用将失败并显示堆栈溢出错误，否则会出现未指定的启动错误。当编译器无法确定堆栈大小时，它会发出警告，指出无法静态确定堆栈大小。递归函数通常就是这种情况。发出此警告后，如果默认堆栈大小不足，用户将需要手动设置堆栈大小。

3.3. 版本控制和兼容性

开发者在开发 CUDA 应用时，应该关注两个版本号：一个是描述计算设备的一般规格和特性的计算能力（详见计算能力），另一个是描述驱动 API 和运行时支持的功能的 CUDA 驱动 API 版本。

驱动程序 API 的版本在驱动程序头文件中定义为。它允许开发人员检查他们的应用程序是否需要比当前安装的设备驱动程序更新的设备驱动程序。这一点很重要，因为驱动程序 API 是向后兼容的，这意味着针对驱动程序 API 的特定版本编译的应用程序、插件和库（包括 CUDA 运行时）将继续在后续设备驱动程序版本上运行，如图 12 所示。驱动程序 API 不向前兼容，这意味着针对特定版本的驱动程序 API 编译的应用程序、插件和库（包括 CUDA 运行时）将无法在以前版本的设备驱动程序上运行。CUDA_VERSION

需要注意的是，支持的版本混合和匹配存在一些限制：

由于系统上一次只能安装一个版本的 CUDA 驱动程序，因此已安装的驱动程序的版本必须与构建必须在该系统上运行的任何应用程序、插件或库所依据的最高驱动程序 API 版本相同或更高。
应用程序使用的所有插件和库都必须使用相同版本的 CUDA 运行时，除非它们静态链接到运行时，在这种情况下，运行时的多个版本可以在同一进程空间中共存。需要注意的是，如果用于链接应用，则默认使用 CUDA Runtime 库的静态版本，所有 CUDA Toolkit 库都静态链接到 CUDA Runtime。nvcc
应用程序使用的所有插件和库必须使用使用运行时的任何库（例如 cuFFT、cuBLAS 等）的相同版本，除非静态链接到这些库。

针对 Tesla GPU 产品，CUDA 10 为 CUDA Driver 的用户模式组件引入了新的向前兼容升级路径。此功能在 CUDA 兼容性中进行了描述。此处所述的 CUDA 驱动程序版本要求适用于用户模式组件的版本。

3.4. 计算模式

在运行 Windows Server 2008 及更高版本或 Linux 的 Tesla 解决方案上，可以使用 NVIDIA 的系统管理接口（nvidia-smi）将系统中的任何设备设置为以下三种模式之一，该接口是作为驱动程序的一部分分发的工具：

默认计算模式：多个主机线程可以同时使用设备（在使用运行时 API 时通过在此设备上调用，或者在使用驱动程序 API 时使当前上下文与设备关联）。cudaSetDevice()
独占进程计算模式：在系统的所有进程中，设备上只能创建一个 CUDA 上下文。在创建该上下文的进程中，上下文可以是任意数量的线程的当前线程。
禁止的计算模式：设备上无法创建 CUDA 上下文。

具体而言，这意味着，如果设备 0 处于禁止模式或独占进程模式并被另一个进程使用，则使用运行时 API 而不显式调用的主机线程可能与设备 0 以外的设备相关联。可用于从按优先级排列的设备列表中设置设备。cudaSetDevice()cudaSetValidDevices()

另请注意，对于采用 Pascal 架构的设备（主要修订号为 6 及更高的计算能力），存在对计算抢占的支持。这使得计算任务可以在指令级粒度上被抢占，而不是像以前的 Maxwell 和 Kepler GPU 架构那样采用线程块粒度，其好处是可以防止具有长时间运行内核的应用程序独占系统或超时。但是，将会产生与计算抢占关联的上下文切换开销，该开销会在支持计算抢占的设备上自动启用。带有该属性的单个属性查询函数可用于确定正在使用的设备是否支持计算抢占。希望避免与不同进程关联的上下文切换开销的用户可以通过选择独占进程模式来确保 GPU 上只有一个进程处于活动状态。cudaDeviceGetAttribute()cudaDevAttrComputePreemptionSupported

应用程序可以通过检查设备属性来查询设备的计算模式（请参阅设备枚举）。computeMode

3.5. 模式开关

具有显示输出的 GPU 将一些 DRAM 内存专用于所谓的主表面，该表面用于刷新用户查看其输出的显示设备。当用户通过更改显示器的分辨率或位深度（使用 NVIDIA 控制面板或 Windows 上的显示器控制面板）来启动显示器的模式切换时，主表面所需的内存量会发生变化。例如，如果用户将显示分辨率从 1280x1024x32 位更改为 1600x1200x32 位，则系统必须将 7.68 MB 专用于主图面，而不是 5.24 MB。（在启用抗锯齿功能的情况下运行的全屏图形应用程序可能需要为主界面提供更多的显示内存。在 Windows 上，可能启动显示模式切换的其他事件包括启动全屏 DirectX 应用程序、按 Alt+Tab 以从全屏 DirectX 应用程序切换任务，或按 Ctrl+Alt+Del 锁定计算机。

如果模式开关增加了主表面所需的内存量，则系统可能不得不蚕食专用于 CUDA 应用程序的内存分配。因此，模式切换会导致对 CUDA 运行时的任何调用都失败并返回无效的上下文错误。

3.6. Windows 的 Tesla 计算集群模式

使用 NVIDIA 的系统管理接口（nvidia-smi），可以将 Windows 设备驱动程序置于 Tesla 和 Quadro 系列设备的 TCC（Tesla 计算集群）模式。

TCC 模式将删除对任何图形功能的支持。

4. 硬件实现

NVIDIA GPU 架构是围绕可扩展的多线程流式多处理器 （SM） 阵列构建的。当主机 CPU 上的 CUDA 程序调用内核网格时，网格的块将被枚举并分发到具有可用执行容量的多处理器。线程块的线程在一个多处理器上并发执行，多个线程块可以在一个多处理器上并发执行。当线程块终止时，将在空出的多处理器上启动新块。

多处理器设计用于同时执行数百个线程。为了管理如此大量的线程，它采用了一种称为 SIMT（单指令多线程）的独特架构，该架构在 SIMT 架构中进行了描述。指令是通过管道传输的，利用单个线程内的指令级并行性，以及通过同时硬件多线程实现的广泛线程级并行性，如硬件多线程中所述。与 CPU 内核不同，它们是按顺序发出的，没有分支预测或推测执行。

SIMT 体系结构和硬件多线程描述了所有设备通用的流式多处理器的体系结构功能。计算能力 5.x、计算能力 6.x 和计算能力 7.x 分别提供计算能力 5.x、6.x 和 7.x 的设备的详细信息。

NVIDIA GPU 架构使用 little-endian 表示。

4.1. SIMT架构

多处理器以 32 个并行线程（称为 warps）为一组创建、管理、调度和执行线程。组成翘曲的各个线程从相同的程序地址一起开始，但它们有自己的指令地址计数器和寄存器状态，因此可以自由地分支和独立执行。“经线”一词起源于编织，这是第一种平行线技术。半经线是经线的前半部分或后半部分。四分之一经线是经线的第一、第二、第三或第四刻度。

当多处理器被赋予一个或多个线程块来执行时，它会将它们划分为扭曲，每个扭曲都由扭曲调度器调度执行。块被划分为扭曲的方式始终是相同的;每个翘曲包含连续的线程，线程 ID 不断增加，第一个翘曲包含线程 0。线程层次结构描述了线程 ID 如何与块中的线程索引相关联。

Warp 一次执行一条公共指令，因此当 Warp 的所有 32 个线程都同意其执行路径时，可以实现全部效率。如果 warp 的线程通过依赖于数据的条件分支发散，则 warp 将执行所采用的每个分支路径，从而禁用不在该路径上的线程。分支发散仅发生在经线内;不同的 Warp 独立执行，无论它们执行的是公共代码路径还是不相交的代码路径。

SIMT 架构类似于 SIMD（单指令多数据）向量组织，因为一条指令控制多个处理元素。一个关键的区别是，SIMD 向量组织向软件公开 SIMD 宽度，而 SIMT 指令指定单个线程的执行和分支行为。与 SIMD 向量机相比，SIMT 使程序员能够为独立的标量线程编写线程级并行代码，并为协调线程编写数据并行代码。为了正确性，程序员基本上可以忽略 SIMT 行为;但是，通过注意代码很少需要 warp 中的线程来发散，可以实现实质性的性能改进。在实践中，这类似于缓存行在传统代码中的作用：在设计正确性时，可以安全地忽略缓存行大小，但在设计峰值性能时，必须在代码结构中考虑这一点。另一方面，向量架构需要软件将负载合并到向量中并手动管理背离。

在 NVIDIA Volta 之前， warps 使用在 warp 中的所有 32 个线程之间共享的单个程序计数器，以及一个指定 warp 的活动线程的活动掩码。因此，来自同一翘曲的线程在不同的区域或不同的执行状态中无法相互发送信号或交换数据，而需要细粒度共享由锁或互斥锁保护的数据的算法很容易导致死锁，具体取决于竞争线程来自哪个翘曲。

从 NVIDIA Volta 架构开始， 独立线程调度允许线程之间完全并发，无论翘曲如何。通过独立线程调度，GPU 维护每个线程的执行状态，包括程序计数器和调用堆栈，并且可以按每个线程的粒度执行，以更好地利用执行资源或允许一个线程等待另一个线程生成数据。计划优化器确定如何将来自同一翘曲的活动线程组合成 SIMT 单元。这与以前的 NVIDIA GPU 一样保留了 SIMT 执行的高吞吐量，但具有更大的灵活性：线程现在可以在亚扭曲粒度上发散和重新收敛。

如果开发人员对扭曲同步性做出假设，则独立线程调度可能会导致一组线程参与执行的代码与预期完全不同2以前的硬件架构。特别是，应重新审视任何 warp-synchronous 代码（例如无同步、warp 内减少），以确保与 NVIDIA Volta 及更高版本兼容。有关详细信息，请参阅计算能力 7.x。

4.2. 硬件多线程

多处理器处理的每个翘曲的执行上下文（程序计数器、寄存器等）在翘曲的整个生命周期内在片上维护。因此，从一个执行上下文切换到另一个执行上下文是没有成本的，并且在每次发出指令时，warp 调度器都会选择一个 warp，该 warp 具有准备好执行其下一条指令的线程（warp 的活动线程），并向这些线程发出指令。

具体而言，每个多处理器都有一组 32 位寄存器，这些寄存器在经线之间分区，还有一个并行数据缓存或共享内存，这些寄存器在线程块之间分区。

对于给定内核，可以在多处理器上驻留和处理的块和扭曲数取决于内核使用的寄存器和共享内存量以及多处理器上可用的寄存器和共享内存量。此外，每个多处理器还有最大驻留块数和最大驻留翘曲数。这些限制以及多处理器上可用的寄存器和共享内存量是器件计算能力的函数，在计算能力中给出。如果每个多处理器没有足够的寄存器或共享内存来处理至少一个块，则内核将无法启动。

一个块中的翘曲总数如下：

$ceil(\frac{T}{W_{size}},1)$

T 是每个块的线程数，
Wsize 是翘曲尺寸，等于 32，
ceil（x， y）等于 x 四舍五入到最接近的 y 的倍数。

寄存器总数和为块分配的共享内存总量记录在 CUDA 工具包中提供的 CUDA 占用计算器中。

术语 warp-synchronous 是指隐含地假设同一 warp 中的线程在每条指令时都同步的代码。

5. 绩效准则

5.1. 整体性能优化策略

性能优化围绕四个基本策略展开：

最大化并行执行，实现最大利用率;
优化内存使用，实现最大内存吞吐量;
优化指令使用，实现最大指令吞吐量;
最大程度地减少内存抖动。

哪些策略将为应用程序的特定部分产生最佳性能增益，取决于该部分的性能限制器;例如，优化主要受内存访问限制的内核的指令使用不会产生任何显著的性能提升。因此，应通过测量和监控性能限制器（例如使用 CUDA 分析器）来不断指导优化工作。此外，将特定内核的浮点运算吞吐量或内存吞吐量（以更有意义的为准）与设备的相应峰值理论吞吐量进行比较，可以表明内核还有多少改进空间。

5.2. 最大化利用率

为了最大限度地提高利用率，应用程序的结构应尽可能公开并行性，并有效地将此并行性映射到系统的各个组件，以使它们大部分时间保持忙碌。

5.2.1. 应用程序级别

概括地说，应用程序应通过使用异步函数调用和流，最大限度地提高主机、设备以及将主机连接到设备的总线之间的并行执行，如异步并发执行中所述。它应该为每个处理器分配它最擅长的工作类型：将工作负载串行到主机;与设备并行的工作负载。

对于并行工作负载，在算法中由于某些线程需要同步以便相互共享数据而导致并行性中断的时刻，有两种情况：要么这些线程属于同一个块，在这种情况下，它们应该在同一内核调用中通过共享内存使用和共享数据，或者它们属于不同的块，在这种情况下，它们必须使用两个单独的内核调用通过全局内存共享数据，一个用于写入，另一个用于从全局内存读取。第二种情况的优化效果要差得多，因为它增加了额外内核调用和全局内存流量的开销。因此，应通过将算法映射到 CUDA 编程模型来最大程度地减少其发生，以便尽可能在单个线程块内执行需要线程间通信的计算。__syncthreads()

5.2.2. 设备级别

在较低级别上，应用程序应最大限度地提高设备的多处理器之间的并行执行。

多个内核可以在一个设备上并发执行，因此也可以通过使用流来实现最大利用率，以使足够的内核能够并发执行，如异步并发执行中所述。

5.2.3. 多处理器级别

在更低的层次上，应用程序应最大限度地提高多处理器内各种功能单元之间的并行执行。

如硬件多线程中所述，GPU 多处理器主要依赖于线程级并行性来最大化其功能单元的利用率。因此，利用率与驻留经线的数量直接相关。在每次发出指令时，warp 调度器都会选择一个准备执行的指令。此指令可以是同一 warp 的另一个独立指令，利用指令级并行性，或者更常见的是另一个 warp 的指令，利用线程级并行性。如果选择了准备执行的指令，则会将其发出给翘曲的活动线程。翘曲准备执行其下一条指令所需的时钟周期数称为延迟，当所有翘曲调度器始终在该延迟期间的每个时钟周期中发出一些指令以处理某些翘曲时，或者换句话说，当延迟完全“隐藏”时，就可以实现充分利用。隐藏 L 个时钟周期的延迟所需的指令数量取决于这些指令的相应吞吐量（有关各种算术指令的吞吐量，请参阅算术指令）。如果我们假设指令具有最大吞吐量，则等于：

4L 适用于计算能力为 5.x、6.1、6.2、7.x 和 8.x 的设备，因为对于这些设备，多处理器在一个时钟周期内每次 warp 发出一条指令，一次四次 warp，如计算能力中所述。
对于计算能力为 6.0 的设备，为 2L，因为对于这些设备，每个周期发出的两条指令是针对两个不同翘曲的一条指令。

warp 尚未准备好执行其下一条指令的最常见原因是指令的输入操作数尚不可用。

如果所有输入操作数都是寄存器，则延迟是由寄存器依赖性引起的，即，某些输入操作数是由尚未执行完成的某些先前指令写入的。在这种情况下，延迟等于前一条指令的执行时间，并且 warp 调度器必须在这段时间内安排其他 warp 的指令。执行时间因指令而异。在计算能力为 7.x 的设备上，对于大多数算术指令，它通常是 4 个时钟周期。这意味着每个多处理器需要 16 个活动 warp（4 个周期，4 个 warp 调度器）来隐藏算术指令延迟（假设 warp 以最大吞吐量执行指令，否则需要更少的 warp）。如果单个 Warper 表现出指令级并行性，即其指令流中有多个独立指令，则需要较少的 Warp，因为来自单个 Warp 的多个独立指令可以背靠背发出。

如果某些输入操作数位于片外存储器中，则延迟要高得多：通常为数百个时钟周期。在如此高的延迟期间，使 warp 调度器保持忙碌所需的 warp 数量取决于内核代码及其指令级并行度。一般来说，如果没有片外存储器操作数的指令数量（即大多数时候是算术指令）与具有片外存储器操作数的指令数量之比较低（这个比率通常称为程序的算术强度），则需要更多的扭曲。

warp 未准备好执行其下一条指令的另一个原因是它正在某个内存围栏（内存围栏函数）或同步点（同步函数）等待。同步点可以强制多处理器空闲，因为越来越多的翘曲等待同一块中的其他翘曲在同步点之前完成指令的执行。在这种情况下，每个多处理器拥有多个驻留块有助于减少空闲，因为来自不同块的翘曲不需要在同步点相互等待。

对于给定的内核调用，驻留在每个多处理器上的块和扭曲的数量取决于调用的执行配置（执行配置）、多处理器的内存资源以及内核的资源要求，如硬件多线程中所述。使用该选项进行编译时，编译器会报告寄存器和共享内存使用情况。--ptxas-options=-v

一个块所需的共享内存总量等于静态分配的共享内存量和动态分配的共享内存量之和。

内核使用的寄存器数量会对驻留扭曲的数量产生重大影响。例如，对于计算能力为 6.x 的设备，如果内核使用 64 寄存器，每个块有 512 个线程，需要很少的共享内存，然后两个块（即 32 个 warp）可以驻留在多处理器上，因为它们需要 2x512x64 寄存器，与多处理器上可用的寄存器数量完全匹配。但是，一旦内核使用另一个寄存器，就只能有一个块（即 16 个 warps）驻留，因为两个块需要 2x512x65 寄存器，这些寄存器比多处理器上可用的寄存器更多。因此，编译器尝试最小化寄存器在保持寄存器溢出（请参阅设备内存访问）和指令数量的同时使用到最低限度。寄存器使用可以是使用编译器选项、Launch Bounds 中所述的限定符或每线程最大寄存器数中所述的限定符进行控制。maxrregcount__launch_bounds__()__maxnreg__()

寄存器文件组织为 32 位寄存器。因此，存储在寄存器中的每个变量至少需要一个 32 位寄存器，例如，一个变量使用两个 32 位寄存器。double

对于给定的内核调用，执行配置对性能的影响通常取决于内核代码。因此，建议进行实验。应用程序还可以根据寄存器文件大小和共享内存大小来参数化执行配置，这取决于设备的计算能力，以及设备的多处理器数量和内存带宽，所有这些都可以使用运行时进行查询（请参阅参考手册）。

每个块的线程数应选择为翘距大小的倍数，以尽可能避免在填充不足的翘距上浪费计算资源。

5.2.3.1. 占用计算器

存在多个 API 函数来帮助程序员根据寄存器和共享内存要求选择线程块大小和集群大小。

占用计算器 API 可以根据内核的块大小和共享内存使用情况提供占用预测。此函数根据每个多处理器的并发线程块数报告占用情况。cudaOccupancyMaxActiveBlocksPerMultiprocessor
- 请注意，此值可以转换为其他指标。乘以每个块的翘曲数，得到每个多处理器的并发翘数;进一步将并发 Warps 除以每个多处理器的最大 Warps 数，得出占用率百分比。
基于占用率的启动配置器 API 和启发式计算实现最大多处理器级别占用的执行配置。cudaOccupancyMaxPotentialBlockSizecudaOccupancyMaxPotentialBlockSizeVariableSMem
占用计算器 API 可以根据内核的集群大小、块大小和共享内存使用情况提供占用预测。此函数根据系统中存在的 GPU 上给定大小的最大活动集群数报告占用情况。cudaOccupancyMaxActiveClusters

以下代码示例计算 MyKernel 的占用率。然后，它报告占用水平以及并发扭曲数与每个多处理器的最大扭曲数之间的比率。

// Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d[idx] = a[idx] * b[idx];
}

// Host code
int main()
{
    int numBlocks;        // Occupancy in terms of active blocks
    int blockSize = 32;

    // These variables are used to convert occupancy to warps
    int device;
    cudaDeviceProp prop;
    int activeWarps;
    int maxWarps;

    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);

    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &numBlocks,
        MyKernel,
        blockSize,
        0);

    activeWarps = numBlocks * blockSize / prop.warpSize;
    maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;

    std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;

    return 0;
}

以下代码示例根据用户输入配置 MyKernel 的基于占用的内核启动。

// Device code
__global__ void MyKernel(int *array, int arrayCount)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arrayCount) {
        array[idx] *= array[idx];
    }
}

// Host code
int launchMyKernel(int *array, int arrayCount)
{
    int blockSize;      // The launch configurator returned block size
    int minGridSize;    // The minimum grid size needed to achieve the
                        // maximum occupancy for a full device
                        // launch
    int gridSize;       // The actual grid size needed, based on input
                        // size

    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize,
        &blockSize,
        (void*)MyKernel,
        0,
        arrayCount);

    // Round up according to array size
    gridSize = (arrayCount + blockSize - 1) / blockSize;

    MyKernel<<<gridSize, blockSize>>>(array, arrayCount);
    cudaDeviceSynchronize();

    // If interested, the occupancy can be calculated with
    // cudaOccupancyMaxActiveBlocksPerMultiprocessor

    return 0;
}

以下代码示例演示如何使用集群占用 API 来查找给定大小的活动集群的最大数量。下面的示例代码计算了大小为 2 的集群的占用率，每个块有 128 个线程。

集群大小为 8 是向前兼容的初始计算能力 9.0，但在 GPU 硬件或 MIG 配置上除外，这些配置太小而无法支持 8 个多处理器，在这种情况下，最大集群大小将减小。但建议用户在启动集群内核之前查询最大集群大小。可以使用 API 查询最大集群大小。cudaOccupancyMaxPotentialClusterSize

{
  cudaLaunchConfig_t config = {0};
  config.gridDim = number_of_blocks;
  config.blockDim = 128; // threads_per_block = 128
  config.dynamicSmemBytes = dynamic_shared_memory_size;

  cudaLaunchAttribute attribute[1];
  attribute[0].id = cudaLaunchAttributeClusterDimension;
  attribute[0].val.clusterDim.x = 2; // cluster_size = 2
  attribute[0].val.clusterDim.y = 1;
  attribute[0].val.clusterDim.z = 1;
  config.attrs = attribute;
  config.numAttrs = 1;

  int max_cluster_size = 0;
  cudaOccupancyMaxPotentialClusterSize(&max_cluster_size, (void *)kernel, &config);

  int max_active_clusters = 0;
  cudaOccupancyMaxActiveClusters(&max_active_clusters, (void *)kernel, &config);

  std::cout << "Max Active Clusters of size 2: " << max_active_clusters << std::endl;
}

CUDA Nsight 计算用户界面还为无法依赖 CUDA 软件堆栈的任何用例提供了一个独立的占用计算器和启动配置器实现。占用计算器的 Nsight Compute 版本作为学习工具特别有用，它可以可视化影响占用的参数（块大小、每个线程的寄存器和每个线程的共享内存）更改的影响。<CUDA_Toolkit_Path>/include/cuda_occupancy.h

5.3. 最大化内存吞吐量

最大化应用程序总体内存吞吐量的第一步是最小化低带宽的数据传输。

这意味着尽量减少主机和设备之间的数据传输，如主机和设备之间的数据传输中所述，因为这些传输的带宽比全局内存和设备之间的数据传输要低得多。

这也意味着通过最大限度地利用片上内存来最大限度地减少全局内存和设备之间的数据传输：共享内存和缓存（即，L1 缓存和 L2 缓存在计算能力为 2.x 及更高的设备上可用，纹理缓存和常量缓存在所有设备上可用）。

共享内存等同于用户管理的缓存：应用程序显式分配并访问它。如 CUDA Runtime 中所示，典型的编程模式是将来自设备内存的数据暂存到共享内存中;换言之，要拥有块的每个线程：

将数据从设备内存加载到共享内存，
与块的所有其他线程同步，以便每个线程都可以安全地读取由不同线程填充的共享内存位置，
在共享内存中处理数据，
如有必要，请再次同步，以确保共享内存已使用结果进行更新。
将结果写回设备内存。

对于某些应用程序（例如，全局内存访问模式依赖于数据），传统的硬件管理缓存更适合利用数据局部性。如计算能力 7.x、计算能力 8.x 和计算能力 9.0 中所述，对于计算能力为 7.x、8.x 和 9.0 的设备，L1 和共享内存使用相同的片上内存，并且每个内核调用都可配置专用于 L1 和共享内存的量。

内核的内存访问吞吐量可能会相差一个数量级，具体取决于每种内存类型的访问模式。因此，最大化内存吞吐量的下一步是根据设备内存访问中所述的最佳内存访问模式，尽可能优化地组织内存访问。这种优化对于全局内存访问尤为重要，因为与可用的片上带宽和算术指令吞吐量相比，全局内存带宽较低，因此非最优全局内存访问通常会对性能产生很大影响。

5.3.1. 主机和设备之间的数据传输

应用程序应努力最大程度地减少主机和设备之间的数据传输。实现此目的的一种方法是将更多代码从主机移动到设备，即使这意味着运行的内核不会暴露出足够的并行性，无法在设备上以全效率执行。中间数据结构可以在设备内存中创建，由设备操作，并在不被主机映射或复制到主机内存的情况下被销毁。

此外，由于每次传输都会产生开销，因此将许多小传输批量转换为单个大传输总是比单独进行每个传输的性能更好。

在具有前端总线的系统上，通过使用页面锁定主机内存，可以实现主机和设备之间的数据传输的更高性能，如页面锁定主机内存中所述。

此外，使用映射页面锁定内存（Mapped Memory）时，无需分配任何设备内存，也无需在设备和主机内存之间显式复制数据。每次内核访问映射内存时，都会隐式执行数据传输。为了获得最佳性能，这些内存访问必须像访问全局内存一样合并（请参阅设备内存访问）。假设它们是，并且映射的内存只被读取或写入一次，那么在设备和主机内存之间使用映射的页面锁定内存而不是显式复制可能会提高性能。

在设备内存和主机内存在物理上相同的集成系统上，主机和设备内存之间的任何副本都是多余的，应改用映射的页面锁定内存。应用程序可以通过检查集成设备属性（请参阅设备枚举）是否等于 1 来查询设备。integrated

5.3.2. 设备内存访问

访问可寻址内存（即全局、本地、共享、常量或纹理内存）的指令可能需要多次重新发出，具体取决于内存地址在 warp 内线程中的分布。这种分布方式如何影响指令吞吐量特定于每种类型的内存，并在以下各节中进行介绍。例如，对于全局内存，作为一般规则，地址越分散，吞吐量就越低。

全局内存

全局内存驻留在设备内存中，设备内存可通过 32、64 或 128 字节内存事务访问。这些内存事务必须自然对齐：内存事务只能读取或写入与其大小对齐的 32、64 或 128 字节的设备内存段（即，其第一个地址是其大小的倍数）。

当 warp 执行访问全局内存的指令时，它会将 warp 内线程的内存访问合并为一个或多个这些内存事务，具体取决于每个线程访问的字的大小以及线程中内存地址的分布。一般来说，需要的事务越多，除了线程访问的单词外，传输的未使用的单词就越多，从而相应地降低了指令吞吐量。例如，如果为每个线程的 4 字节访问生成 32 字节的内存事务，则吞吐量将除以 8。

需要多少事务以及最终影响多少吞吐量取决于设备的计算能力。计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 提供了有关如何处理各种计算功能的全局内存访问的更多详细信息。

因此，为了最大限度地提高全局内存吞吐量，必须通过以下方式最大化合并：

遵循基于计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 的最佳访问模式
使用满足以下“大小和对齐要求”部分中详述的大小和对齐要求的数据类型，
在某些情况下填充数据，例如，在访问二维数组时，如下面的二维数组部分所述。

尺寸和对齐要求

全局内存指令支持读取或写入大小等于 1、2、4、8 或 16 字节的字。当且仅当数据类型的大小为 1、2、4、8 或 16 个字节，并且数据自然对齐（即其地址是该大小的倍数）时，对驻留在全局内存中的数据的任何访问（通过变量或指针）都会编译为单个全局内存指令。

如果不满足此大小和对齐要求，则访问将编译为具有交错访问模式的多个指令，从而阻止这些指令完全合并。因此，对于驻留在全局内存中的数据，建议使用满足此要求的类型。

对于内置向量类型，将自动满足对齐要求。

对于结构，编译器可以使用对齐说明符来强制执行大小和对齐要求，例如__align__(8) or __align__(16)

struct __align__(8) {
    float x;
    float y;
};

或

struct __align__(16) {
    float x;
    float y;
    float z;
};

驻留在全局内存中或由驱动程序或运行时 API 的内存分配例程之一返回的变量的任何地址始终与至少 256 个字节对齐。

读取未自然对齐的 8 字节或 16 字节单词会产生错误的结果（相差几个字），因此必须特别注意保持这些类型的任何值或值数组的起始地址的对齐。一个可能很容易被忽视的典型情况是，在使用一些自定义的全局内存分配方案时，其中多个数组的分配（多次调用或）被分配到多个数组中的单个大内存块所取代，在这种情况下，每个数组的起始地址与块的起始地址相偏移。cudaMalloc()cuMemAlloc()

二维数组

一种常见的全局内存访问模式是，当索引的每个线程使用以下地址来访问宽度为 2D 数组的一个元素时，该数组位于类型的地址（满足最大化利用率中描述的要求）：(tx,ty)widthBaseAddresstype*type

BaseAddress + width * ty + tx

要使这些访问完全合并，线程块的宽度和数组的宽度都必须是翘距大小的倍数。

具体来说，这意味着，如果实际上分配的数组的宽度四舍五入到此大小的最接近倍数，并且相应地填充其行，那么该数组的宽度将得到更有效的访问。参考手册中描述的 and 函数和关联的内存复制函数使程序员能够编写不依赖于硬件的代码来分配符合这些约束的数组。cudaMallocPitch()cuMemAllocPitch()

本地内存

仅对变量内存空间说明符中提到的某些自动变量进行本地内存访问。编译器可能放置在本地内存中的自动变量包括：

它无法确定它们是否以常量进行索引的数组，
会占用太多寄存器空间的大型结构或数组，
如果内核使用的寄存器多于可用寄存器，则为任何变量（这也称为寄存器溢出）。

检查 PTX 汇编代码（通过使用 or 选项编译获得）将判断变量是否在第一个编译阶段已放置在本地内存中，因为它将使用助记符声明并使用和助记符进行访问。即使没有，后续的编译阶段仍可能做出其他决定，但如果他们发现它为目标架构占用了太多的寄存器空间：检查使用的 cubin 对象将判断是否是这种情况。此外，在使用该选项进行编译时，编译器会报告每个内核（）的总本地内存使用量。请注意，某些数学函数具有可能访问本地内存的实现路径。-ptx-keep.localld.localst.localcuobjdumplmem--ptxas-options=-v

本地内存空间驻留在设备内存中，因此本地内存访问具有与全局内存访问相同的高延迟和低带宽，并且受设备内存访问中所述的相同内存合并要求的约束。但是，本地内存的组织方式使得连续的 32 位字可以由连续的线程 ID 访问。因此，只要 warp 中的所有线程都访问相同的相对地址（例如，数组变量中的相同索引，结构变量中的相同成员），访问就完全合并。

在计算能力 5.x 及更高版本的设备上，本地内存访问始终以与全局内存访问相同的方式缓存在 L2 中（请参阅计算能力 5.x 和计算能力 6.x）。

共享内存

因为它是片上的，所以共享内存比本地或全局内存具有更高的带宽和低得多的延迟。

为了实现高带宽，共享内存被划分为大小相等的内存模块，称为存储区，可以同时访问。因此，由位于 n 个不同内存组中的 n 个地址发出的任何内存读取或写入请求都可以同时得到服务，从而产生比单个模块带宽高 n 倍的总带宽。

但是，如果内存请求的两个地址位于同一内存组中，则存在内存组冲突，必须序列化访问。硬件会根据需要将具有银行冲突的内存请求拆分为尽可能多的独立无冲突请求，从而将吞吐量降低一个因子，该系数等于单独内存请求的数量。如果单独的内存请求数为 n，则表示初始内存请求会导致 n 向组冲突。

因此，为了获得最大性能，了解内存地址如何映射到内存组非常重要，以便安排内存请求，从而最大程度地减少内存组冲突。计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 分别在计算能力 5.x、计算能力 6.x、计算能力 8.x、计算能力 9.0 中对此进行了介绍。

恒定内存

常量内存空间驻留在设备内存中，并缓存在常量缓存中。

然后，一个请求被拆分为与初始请求中具有不同内存地址的任意数量的单独请求，从而将吞吐量降低一个因子，该系数等于单独请求的数量。

然后，如果发生缓存命中，则按常量缓存的吞吐量提供服务，否则以设备内存的吞吐量为服务。

纹理和表面记忆

纹理和表面内存空间驻留在设备内存中，并缓存在纹理缓存中，因此纹理提取或表面读取仅在缓存未命中时从设备内存中读取一次内存，否则只需从纹理缓存中读取一次。纹理缓存针对 2D 空间局部性进行了优化，因此，读取在 2D 中彼此靠近的纹理或表面地址的相同翘曲线程将获得最佳性能。此外，它还设计用于具有恒定延迟的流式提取;缓存命中会减少 DRAM 带宽需求，但不会降低提取延迟。

通过纹理或表面提取读取设备内存具有一些优势，可以使其成为从全局或恒定内存中读取设备内存的有利替代方法：

如果内存读取不遵循全局或常量内存读取必须遵循的访问模式以获得良好的性能，则可以实现更高的带宽，前提是纹理提取或表面读取中存在局部性;
寻址计算由专用单元在内核外部执行;
打包的数据可以在单个操作中广播到单独的变量;
8 位和 16 位整数输入数据可以选择性地转换为 [0.0， 1.0] 或 [-1.0， 1.0] 范围内的 32 位浮点值（请参阅纹理内存）。

5.4. 最大化指令吞吐量

为了最大限度地提高指令吞吐量，应用程序应：

尽量减少使用低吞吐量的算术指令;这包括在不影响最终结果的情况下用精度换取速度，例如使用内部函数而不是常规函数（内部函数列在内部函数中）、单精度而不是双精度，或将非规范化数字刷新为零;
最大程度地减少由控制流指令引起的发散翘曲，如控制流指令中所述
减少指令的数量，例如，尽可能优化同步点（如同步指令中所述）或使用受限制的指针（如__restrict__中所述）。

在本节中，吞吐量以每个多处理器每个时钟周期的操作数表示。对于 32 的 warp 大小，一条指令对应于 32 个操作，因此，如果 N 是每个时钟周期的操作数，则指令吞吐量为 N/32 条/每个时钟周期的指令。

所有吞吐量都适用于一个多处理器。它们必须乘以设备中的多处理器数量，才能获得整个设备的吞吐量。

6. 支持 CUDA 的 GPU

CUDA GPUs - Compute Capability | NVIDIA Developer 列出了所有启用了 CUDA 的设备及其计算能力。

可以使用运行时查询计算能力、多处理器数量、时钟频率、设备内存总量和其他属性（请参阅参考手册）。

7. C++ 语言扩展

7.1. 函数执行空间说明符

函数执行空间说明符表示函数是在主机上执行还是在设备上执行，以及它是可从主机还是从设备调用。

7.1.1. global

执行空间说明符将函数声明为内核。这样的函数是：__global__

在设备上执行，
可从主机调用，
对于计算能力为 5.0 或更高版本的设备，可从设备调用（有关详细信息，请参阅 CUDA 动态并行性）。

函数必须具有 void 返回类型，并且不能是类的成员。__global__

对函数的任何调用都必须指定其执行配置，如执行配置中所述。__global__

对函数的调用是异步的，这意味着它在设备完成执行之前返回。__global__

7.1.2. device

执行空间说明符声明了一个函数，该函数为：__device__

在设备上执行，
只能从设备调用。

和执行空间说明符不能一起使用。__global____device__

7.1.3. host

执行空间说明符声明了一个函数，该函数为：__host__

在主机上执行，
只能从主机调用。

它等效于声明一个仅带有执行空间说明符的函数，或者声明一个不带任何、或执行空间说明符的函数;在任何一种情况下，该函数都仅针对主机进行编译。__host____host____device____global__

和执行空间说明符不能一起使用。__global____host__

但是，和执行空间说明符可以一起使用，在这种情况下，将针对主机和设备编译函数。应用程序兼容性中引入的巨集可用于区分主机和设备之间的代码路径：__device____host____CUDA_ARCH__

__host__ __device__ func()
{
#if __CUDA_ARCH__ >= 800
   // Device code path for compute capability 8.x
#elif __CUDA_ARCH__ >= 700
   // Device code path for compute capability 7.x
#elif __CUDA_ARCH__ >= 600
   // Device code path for compute capability 6.x
#elif __CUDA_ARCH__ >= 500
   // Device code path for compute capability 5.x
#elif !defined(__CUDA_ARCH__)
   // Host code path
#endif
}

7.1.4. 未定义的行为

在以下情况下，“交叉执行空间”调用具有未定义的行为：

__CUDA_ARCH__定义，从或函数内部对函数的调用。__global____device____host__ __device____host__
__CUDA_ARCH__是 undefined，即从函数内部到函数的调用。__host____device__9

7.1.5. noinline和forceinline

编译器在认为适当时内联任何函数。__device__

如果可能，函数限定符可以用作编译器不内联函数的提示。__noinline__

函数限定符可用于强制编译器内联函数。__forceinline__

and 函数限定符不能一起使用，并且两个函数限定符都不能应用于内联函数。__noinline____forceinline__

7.1.6. __inline_hint__

限定符在编译器中启用更积极的内联。与不同，它并不意味着函数是内联的。在使用 LTO 时，它可用于改进模块之间的内联。__inline_hint____forceinline__

the 和 function 限定符都不能与 function 限定符一起使用。__noinline____forceinline____inline_hint__

7.2. 可变内存空间说明符

变量内存空间说明符表示变量在设备上的内存位置。

在设备代码中声明的自动变量，不带本节中描述的任何、和内存空间说明符，通常位于寄存器中。但是，在某些情况下，编译器可能会选择将其放置在本地内存中，这可能会产生不利的性能后果，如设备内存访问中所述。__device____shared____constant__

7.2.1. device

内存空间说明符声明驻留在设备上的变量。__device__

最多可以与以下三个部分中定义的其他内存空间说明符中的一个一起使用，以进一步表示变量属于哪个内存空间。如果它们都不存在，则变量：__device__

驻留在全局内存空间中，
具有创建它的 CUDA 上下文的生命周期，
每个设备都有一个不同的对象，
可从网格内的所有线程访问，并通过运行时库从主机访问 / / ）。(cudaGetSymbolAddress()cudaGetSymbolSize()cudaMemcpyToSymbol()cudaMemcpyFromSymbol()

7.2.2. constant

内存空间说明符（可选地与一起使用）声明了一个变量，该变量：__constant____device__

驻留在恒定的内存空间中，
具有创建它的 CUDA 上下文的生命周期，
每个设备都有一个不同的对象，
可从网格内的所有线程访问，也可以通过运行时库（ / / / ）从主机访问。cudaGetSymbolAddress()cudaGetSymbolSize()cudaMemcpyToSymbol()cudaMemcpyFromSymbol()

当存在一个并发网格在此网格生命周期的任何点访问该常量时，从主机修改常量的行为是不确定的。

7.2.3. shared

内存空间说明符（可选地与一起使用）声明了一个变量，该变量：__shared____device__

驻留在线程块的共享内存空间中，
具有块的生命周期，
每个块都有一个不同的对象，
只能从块内的所有线程访问，
没有固定的地址。

将共享内存中的变量声明为外部数组时，例如

extern __shared__ float shared[];

数组的大小是在启动时确定的（请参阅执行配置）。以这种方式声明的所有变量都从内存中的同一地址开始，因此必须通过偏移量显式管理数组中变量的布局。例如，如果想要等效的

short array0[128];
float array1[64];
int   array2[256];

在动态分配的共享内存中，可以按以下方式声明和初始化数组：

extern __shared__ float array[];
__device__ void func()      // __device__ or __global__ function
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[128];
    int*   array2 =   (int*)&array1[64];
}

请注意，指针需要与它们指向的类型对齐，因此，例如，以下代码不起作用，因为 array1 未对齐到 4 个字节。

extern __shared__ float array[];
__device__ void func()      // __device__ or __global__ function
{
    short* array0 = (short*)array;
    float* array1 = (float*)&array0[127];
}

7.2.4. __grid_constant__

大于或等于 7.0 的计算体系结构的注释对非引用类型的限定函数参数进行注释，该参数：__grid_constant__const__global__

具有电网的寿命，
对网格是私有的，即，主机线程和来自其他网格（包括子网格）的线程无法访问该对象，
每个网格都有一个不同的对象，即网格中的所有线程都看到相同的地址，
是只读的，即修改对象或其任何子对象是未定义的行为，包括成员。__grid_constant__mutable

要求：

带注释的内核参数必须具有限定的非引用类型。__grid_constant__const
所有函数声明必须与任何参数匹配。__grid_constant_
函数模板专用化必须与任何参数的主模板声明匹配。__grid_constant__
函数模板实例化指令必须与任何参数的主模板声明匹配。__grid_constant__

如果获取函数参数的地址，编译器通常会在线程本地内存中复制内核参数，并使用副本的地址，以部分支持 C++ 语义，该语义允许每个线程修改自己的函数参数本地副本。对函数参数进行注释可确保编译器不会在线程本地内存中创建内核参数的副本，而是使用参数本身的通用地址。避免本地复制可能会提高性能。__global____global____grid_constant__

__device__ void unknown_function(S const&);
__global__ void kernel(const __grid_constant__ S s) {
   s.x += threadIdx.x;  // Undefined Behavior: tried to modify read-only memory

   // Compiler will _not_ create a per-thread thread local copy of "s":
   unknown_function(s);
}

7.2.5. managed

内存空间说明符（可选地与一起使用）声明了一个变量，该变量：__managed____device__

可以从设备和主机代码中引用，例如，可以获取其地址，也可以直接从设备或主机函数中读取或写入。
具有应用程序的生存期。

有关详细信息，请参阅 __managed__ 内存空间说明符。

7.2.6. restrict

nvcc支持通过关键字限制指针。__restrict__

C99 中引入了受限制指针，以缓解 C 类型语言中存在的混叠问题，并抑制了从代码重新排序到常见子表达式消除的所有类型的优化。

下面是一个存在混叠问题的示例，其中使用受限指针可以帮助编译器减少指令数量：

void foo(const float* a,
         const float* b,
         float* c)
{
    c[0] = a[0] * b[0];
    c[1] = a[0] * b[0];
    c[2] = a[0] * b[0] * a[1];
    c[3] = a[0] * a[1];
    c[4] = a[0] * b[0];
    c[5] = b[0];
    ...
}

在 C 类语言中，指针、和可以被别名，因此任何写入都可能修改或的元素。这意味着，为了保证功能的正确性，编译器不能将和加载到寄存器中，将它们相乘，并将结果存储到和，因为如果结果与抽象执行模型不同，例如，如果与的位置完全相同。因此，编译器无法利用通用子表达式。同样，编译器不能只是将计算的重新排序为计算的邻近性，因为前面的写入可能会改变对计算的输入。abccaba[0]b[0]c[0]c[1]a[0]c[0]c[4]c[0]c[1]c[3]c[4]

通过创建、和 restricted 指针，程序员向编译器断言指针实际上没有被别名，在这种情况下，这意味着写入永远不会覆盖或的元素。这将更改函数原型，如下所示：abccab

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c);

请注意，所有指针参数都需要受到限制，编译器优化器才能获得任何好处。添加关键字后，编译器现在可以随意重新排序和消除常见的子表达式，同时保留与抽象执行模型相同的功能：__restrict__

void foo(const float* __restrict__ a,
         const float* __restrict__ b,
         float* __restrict__ c)
{
    float t0 = a[0];
    float t1 = b[0];
    float t2 = t0 * t1;
    float t3 = a[1];
    c[0] = t2;
    c[1] = t2;
    c[4] = t2;
    c[2] = t2 * t3;
    c[3] = t0 * t3;
    c[5] = t1;
    ...
}

这里的效果是减少了内存访问次数和计算次数。这与由于“缓存”负载和常见子表达式导致的寄存器压力增加相平衡。

由于寄存器压力是许多 CUDA 代码中的一个关键问题，因此由于占用率降低，使用受限指针可能会对 CUDA 代码产生负面性能影响。

7.36. 动态全局内存分配和操作

只有计算能力为 2.x 及更高版本的设备才支持动态全局内存分配和操作。

__host__ __device__ void* malloc(size_t size);
__device__ void *__nv_aligned_device_malloc(size_t size, size_t align);
__host__ __device__  void free(void* ptr);

从全局内存中固定大小的堆中动态分配和释放内存。

__host__ __device__ void* memcpy(void* dest, const void* src, size_t size);

将字节从所指向的内存位置复制到所指向的内存位置。sizesrcdest

__host__ __device__ void* memset(void* ptr, int value, size_t size);

设置 to 指向的内存块的字节数（解释为无符号字符）。sizeptrvalue

CUDA 内核函数至少从设备堆中分配字节，并返回指向已分配内存的指针，如果内存不足，则返回 NULL，如果内存不足，则返回 NULL。返回的指针保证与 16 字节边界对齐。malloc()size

CUDA 内核函数至少从设备堆中分配字节，并返回指向已分配内存的指针，如果内存不足，则返回 NULL 以满足请求的大小或对齐方式。分配的内存的地址将是的倍数。必须是 2 的非零次幂。__nv_aligned_device_malloc()sizealignalign

CUDA 内核内函数解除分配指向的内存，该内存必须由先前对 or 的调用返回。如果为 NULL，则忽略对的调用。重复调用相同的内容具有未定义的行为。free()ptrmalloc()__nv_aligned_device_malloc()ptrfree()free()ptr

给定 CUDA 线程通过 CUDA 上下文分配的内存，或在 CUDA 上下文的生命周期内保持分配的内存，或者直到通过调用显式释放该内存。它可以被任何其他 CUDA 线程使用，甚至可以在后续的内核启动中使用。任何 CUDA 线程都可以释放另一个线程分配的内存，但应注意确保不会多次释放同一指针。malloc()__nv_aligned_device_malloc()free()

7.36.3. 示例

7.36.3.1. 每线程分配

以下代码示例：

#include <stdlib.h>
#include <stdio.h>

__global__ void mallocTest()
{
    size_t size = 123;
    char* ptr = (char*)malloc(size);
    memset(ptr, 0, size);
    printf("Thread %d got pointer: %p\n", threadIdx.x, ptr);
    free(ptr);
}

int main()
{
    // Set a heap size of 128 megabytes. Note that this must
    // be done before any kernel is launched.
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
    mallocTest<<<1, 5>>>();
    cudaDeviceSynchronize();
    return 0;
}

输出：

Thread 0 got pointer: 0x20f9ffe20
Thread 1 got pointer: 0x20f9ffec0
Thread 2 got pointer: 0x20f9fff60
Thread 3 got pointer: 0x20f9f97c0
Thread 4 got pointer: 0x20f9f9720