CUDA C++为熟悉C++编程语言的用户提供了一个简单的路径，使他们能够轻松地编写供设备执行的程序。
它由C++语言的最小扩展集和一个运行时库组成。
核心语言扩展已在编程模型中介绍。它们允许程序员将内核定义为C++函数，并使用一些新的语法每次调用函数时指定grid and block的维度。所有扩展的完整描述可以在C++语言扩展中找到。包含这些扩展的任何源文件都必须按照NVCC编译中的说明使用nvcc进行编译。
运行时在CUDA运行时中被介绍。它提供在主机上执行的C和C++函数，用于分配和释放设备内存，传输主机内存和设备内存之间的数据，管理具有多个设备的系统等。运行时的完整描述可以在CUDA参考手册中找到。
运行时是基于一个更低级别的C API构建的，即CUDA驱动API，应用程序也可以访问。驱动API通过公开更低级别的概念（如CUDA上下文 - 设备的主机进程的类比，以及CUDA模块 - 设备的动态加载库的类比）提供了额外的控制级别。大多数应用程序不使用驱动API，因为它们不需要这个额外的控制级别，而当使用运行时时，上下文和模块管理是隐式的，从而产生更简洁的代码。由于运行时与驱动API是互操作的，大多数需要一些驱动API功能的应用程序可以默认使用运行时API，并仅在需要时使用驱动API。驱动程序API在Driver API中介绍。
使用NVCC进行编译
内核可以使用CUDA指令集架构PTX编写，该架构在PTX参考手册中有描述。然而，通常更有效的方法是使用高级编程语言如C++。在两种情况下，内核必须由nvcc编译成二进制代码才能在设备上执行。
nvcc是一个编译器驱动程序，简化了编译C++或PTX代码的过程：它提供简单和熟悉的命令行选项，并通过调用实现不同编译阶段的工具集合来执行它们。本节概述了nvcc工作流程和命令选项。完整说明可在nvcc用户手册中找到。
Compilation Workflow
Offline Compilation 离线编译
使用nvcc编译的源文件可以包括host代码（即在主机上执行的代码）和设备代码（即在设备上执行的代码）的混合。nvcc的基本工作流程包括将设备代码与主机代码分开，然后:
- 将设备代码编译成汇编形式（PTX代码）和/或二进制形式（cubin对象）,
- 并通过替换在内核中引入的<<<…>>>语法（并在执行配置中更详细地描述）来修改主机代码，以必要的CUDA运行时函数调用从PTX代码和/或cubin对象加载和启动每个编译的内核
修改后的主机代码输出为C++代码，可以使用另一个工具编译，或者直接通过让nvcc在最后的编译阶段调用主机编译器输出为对象代码。
Applications can then:
- 要么链接到编译后的主机代码（这是最常见的情况），
- 要么忽略修改后的主机代码（如果有的话），并使用CUDA驱动API（参见驱动API）加载和执行PTX代码或cubin对象。
Just-in-Time Compilation 即时编译
任何在运行时由应用程序加载的PTX代码都会被设备驱动程序进一步编译成二进制代码。这被称为即时编译。即时编译增加了应用程序的加载时间，但允许应用程序受益于每个新设备驱动程序带来的任何新的编译器改进。这也是应用程序在被编译时尚不存在的设备上运行的唯一方式，如在应用程序兼容性中详细描述。
当设备驱动程序为某个应用程序即时编译一些PTX代码时，它会自动缓存生成的二进制代码的副本，以避免在后续调用应用程序时重复编译。该缓存 - 称为计算缓存 - 在升级设备驱动程序时会自动失效，这样应用程序就可以从新的即时编译器中受益，该编译器内置于设备驱动程序中。
环境变量可用于控制即时编译，如在CUDA环境变量中所述。
作为使用nvcc编译CUDA C++设备代码的替代方法，可以在运行时使用NVRTC编译CUDA C++设备代码到PTX。NVRTC是一个用于CUDA C++的运行时编译库；更多信息可以在NVRTC用户指南中找到。
Binary Compatibility
二进制代码是特定于架构的。使用编译器选项-code 生成cubin 对象，该选项指定目标架构：例如，使用-code=sm_80 编译会为计算能力 8.0的设备生成二进制代码。从一个次要版本到下一个次要版本保证二进制兼容性，但从一个次要版本到前一个次要版本或跨主要版本则不保证。换句话说，为计算能力X.y 生成的cubin 对象只会在计算能力为X.z 的设备上执行，其中z≥y 。
Binary compatibility is supported only for the desktop. It is not supported for Tegra. Also, the binary compatibility between desktop and Tegra is not supported.
PTX Compatibility
某些PTX指令仅在具有更高计算能力的设备上受支持。例如，Warp Shuffle函数仅在计算能力为5.0及以上的设备上受支持。-arch编译器选项指定编译C++到PTX代码时假定的计算能力。因此，包含warp shuffle的代码，例如，必须使用-arch=compute_50（或更高）进行编译。
为某个特定计算能力生成的PTX代码始终可以编译为更大或相等的计算能力的二进制代码。请注意，从早期PTX版本编译的二进制文件可能不会使用某些硬件功能。例如，针对计算能力7.0（Volta）的设备编译的二进制文件，从为计算能力6.0（Pascal）生成的PTX编译，将不会使用Tensor Core指令，因为Pascal上没有这些指令。因此，最终的二进制文件的性能可能比使用PTX的最新版本生成的二进制文件差。
Application Compatibility
为了在具有特定计算能力的设备上执行代码，应用程序必须加载与此计算能力兼容的二进制或PTX代码，如在二进制兼容性和PTX兼容性中所述。特别地，为了能够在具有更高计算能力的未来架构上执行代码（为其尚未生成二进制代码），应用程序必须加载将为这些设备即时编译的PTX代码（参见即时编译）。
CUDA C++应用程序中嵌入的PTX和二进制代码是由-arch和-code编译器选项或-gencode编译器选项控制的，如nvcc用户手册中详细描述的那样。例如，
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兼容的二进制代码（第一和第二个-gencode选项）以及与计算能力7.0兼容的PTX和二进制代码（第三个-gencode选项）。
主机代码在运行时自动生成，以自动选择最适合的代码进行加载和执行，上述示例中将是：
- 5.0的二进制代码用于计算能力为5.0和5.2的设备，
- 6.0的二进制代码用于计算能力为6.0和6.1的设备，
- 7.0的二进制代码用于计算能力为7.0和7.5的设备，
- PTX代码在运行时编译为计算能力为8.0和8.6的设备的二进制代码。
x.cu可以有一个优化的代码路径，例如使用仅在计算能力为8.0及更高版本的设备上受支持的warp reduction操作。__CUDA_ARCH__宏可用于基于计算能力区分各种代码路径。它仅为设备代码定义。例如，当使用-arch=compute_80编译时，__CUDA_ARCH__等于800。
使用驱动API的应用程序必须将代码编译为单独的文件，并在运行时明确加载和执行最合适的文件。
Volta架构引入了独立线程调度，这改变了GPU上线程的调度方式。对于依赖于先前架构中SIMT调度的特定行为的代码，独立线程调度可能会更改参与线程的集合，导致结果不正确。为了在实施独立线程调度中详细描述的纠正措施时帮助迁移，Volta开发人员可以选择使用编译器选项组合-arch=compute_60 -code=sm_70来选择Pascal的线程调度。
nvcc用户手册列出了-arch、-code和-gencode编译器选项的各种简写。例如，-arch=sm_70是-arch=compute_70 -code=compute_70,sm_70的简写（与-gencode arch=compute_70,code="compute_70,sm_70"相同）。
C++ Compatibility
编译器的前端根据C++语法规则处理CUDA源文件。主机代码支持完整的C++。但是，只有C++语言支持中描述的部分C++代码在device代码中得到完全支持。
64-位兼容
nvcc的64位版本以64位模式编译设备代码（即指针为64位）。以64位模式编译的设备代码只能与以64位模式编译的主机代码一起使用。
CUDA Runtime
运行时在cudart库中实现，该库链接到应用程序，可以通过cudart.lib或libcudart.a静态链接，或通过cudart.dll或libcudart.so动态链接。需要cudart.dll和/或cudart.so进行动态链接的应用程序通常将它们作为应用程序安装包的一部分包含在内。只有在链接到CUDA运行时的同一实例的组件之间传递CUDA运行时符号的地址才是安全的。
它的所有入口点都带有cuda前缀。
- 如异构编程中所述，CUDA编程模型假定一个由主机和设备组成的系统，每个系统都有自己的独立内存。设备内存概述了用于管理设备内存的运行时函数。
- 共享内存说明了如何使用在线程层次结构中引入的共享内存来最大化性能。
- 页锁定的主机内存介绍了与设备内存之间的数据传输重叠的内核执行所需的页锁定的主机内存。
- 异步并发执行描述了用于在系统的各个级别启用异步并发执行的概念和API。
- 多设备系统展示了编程模型如何扩展到与同一主机连接的多个设备的系统。
- 错误检查描述了如何正确检查运行时生成的错误。
- 调用堆栈提到了用于管理CUDA C++调用堆栈的运行时函数。
- 纹理和表面内存介绍了纹理和表面内存空间，这些空间提供了另一种访问设备内存的方法；它们还公开了GPU纹理硬件的一个子集。
- 图形互操作性介绍了运行时提供的与两个主要图形API（OpenGL和Direct3D）互操作的各种函数。
Initialization
从CUDA 12.0开始，cudaInitDevice()和cudaSetDevice()调用会初始化与指定设备关联的运行时和主上下文。如果没有这些调用，运行时将隐式使用设备0，并根据需要自我初始化以处理其他运行时API请求。在计时运行时函数调用和解释第一次调用运行时的错误代码时，需要记住这一点。在12.0之前，cudaSetDevice()不会初始化运行时，应用程序通常会使用无操作运行时调用cudaFree(0)来隔离运行时初始化和其他api活动（为了计时和错误处理）。
运行时为系统中的每个设备创建一个CUDA上下文（有关CUDA上下文的更多详细信息，请参见上下文）。这个上下文是该设备的主上下文，并在第一个需要在此设备上有活动上下文的运行时函数上初始化。它在应用程序的所有主机线程之间共享。作为此上下文创建的一部分，如果需要，设备代码会进行即时编译（参见即时编译）并加载到设备内存中。这一切都是透明的。如果需要，例如，为了驱动API互操作性，可以从驱动API访问设备的主上下文，如运行时和驱动API之间的互操作性中所述。
当主机线程调用cudaDeviceReset()时，这将销毁主机线程当前操作的设备的主上下文（即，在设备选择中定义的当前设备）。任何将此设备作为当前设备的主机线程进行的下一个运行时函数调用都将为此设备创建一个新的主上下文。
CUDA接口使用在主机程序初始化期间初始化并在主机程序终止期间销毁的全局状态。CUDA运行时和驱动程序无法检测此状态是否无效，因此在程序初始化或终止期间（在main之后）使用这些接口（隐式或显式）将导致未定义的行为。
从CUDA 12.0开始，cudaSetDevice()现在将在为主机线程更改当前设备后明确初始化运行时。CUDA的先前版本在cudaSetDevice()之后延迟了新设备上的运行时初始化，直到进行了第一个运行时调用。这一变化意味着现在检查cudaSetDevice()的返回值以查找初始化错误非常重要。
参考手册的错误处理和版本管理部分的运行时函数不会初始化运行时。
Device Memory
如在异构编程中所提及的，CUDA编程模型假设一个由主机和设备组成的系统，每个都有自己的独立内存。内核在设备内存中运行，因此运行时提供了分配、释放和复制设备内存的函数，以及在主机内存和设备内存之间传输数据。
设备内存可以分配为线性内存或CUDA数组。
CUDA数组是为纹理提取优化的不透明内存布局。它们在 Texture and Surface Memory中有描述。
线性内存在单一统一的地址空间中分配，这意味着单独分配的实体可以通过指针相互引用，例如在二叉树或链表中。地址空间的大小取决于主机系统（CPU）和所使用的GPU的计算能力：
在计算能力为5.3（Maxwell）及更早版本的设备上，CUDA驱动程序会创建一个未提交的40位虚拟地址保留，以确保内存分配（指针）位于支持的范围内。这个保留会显示为保留的虚拟内存，但直到程序实际分配内存之前，不会占用任何物理内存。
线性内存通常使用cudaMalloc()进行分配，并使用cudaFree()进行释放，主机内存和设备内存之间的数据传输通常使用cudaMemcpy()完成。在内核的向量加法代码示例中，需要将向量从主机内存复制到设备内存：
// Device code
__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];
}
// Host code
int main()
{
int N = ...;
size_t size = N * sizeof(float);
// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
float* h_C = (float*)malloc(size);
// Initialize input vectors
...
// 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
// h_C contains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
// Free host memory
...
}