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 的