5.3 内存层次 Memory Hierarchy
CUDA线程在执行过程中可能会访问多个内存空间的数据，如图6所示。每个线程都有自己的私有本地内存。
每个线程块都有一个对块内所有线程可见的共享内存，并且其生命周期与块相同。线程块集群中的线程块可以对彼此的共享内存执行读、写和原子操作。所有线程都可以访问同一块全局内存。
此外，还有两个只读内存空间可以被所有线程访问：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间都针对不同的内存使用进行了优化（参见设备内存访问章节）。纹理内存也提供了不同的寻址模式，以及针对某些特定数据格式的数据过滤（参见纹理和表面内存章节）。
全局内存(global)、常量内存(constant)和纹理内存(texture)空间在相同应用程序的内核启动间都是持久的。
图6：内存层次
5.4 异构编程
如图7所示，CUDA编程模型假设CUDA线程在一个物理独立的设备上运行，该设备作为运行C++程序的主机的协处理器。例如，当内核在GPU上运行，而C++程序的其余部分在CPU上运行时，就是这种情况。
CUDA编程模型还假设主机和设备在DRAM中分别维护自己的独立内存空间，分别称为主机内存和设备内存。因此，一个程序通过调用CUDA运行时（在编程接口章节中描述）来管理内核可以看到的全局内存、常量内存和纹理内存空间。这包括设备内存的分配和释放，以及主机和设备内存之间的数据传输。
统一内存提供了管理内存，以连接主机和设备的内存空间。管理内存可以作为一个统一、连贯的内存映像，通过一个共享的地址空间，从系统中的所有CPU和GPU访问。这一能力使设备内存能被过度订阅，并且可以大大简化转换应用程序的任务，因为它消除了在主机和设备之间明确镜像数据的需要。请参阅统一内存编程章节来了解统一内存的介绍。
图7：异构编程
串行代码在主机上执行，而并行代码在设备上执行
5.5 异步SIMT编程模型
在CUDA编程模型中，线程是执行计算或内存操作的最低级别的抽象。从基于NVIDIA Ampere GPU架构的设备开始，CUDA编程模型通过异步编程模型为内存操作提供加速。异步编程模型定义了异步操作相对于CUDA线程的行为。
异步编程模型定义了异步屏障的行为，用于CUDA线程之间的同步。该模型还解释和定义了如何使用cuda::memcpy_async在GPU进行计算的同时异步地从全局内存移动数据。
5.5.1 异步操作
异步操作被定义为由CUDA线程启动并由另一个线程异步执行的操作。在一个规范的程序中，一个或多个CUDA线程与异步操作同步。启动异步操作的CUDA线程并不需要在同步线程中。
这样的异步线程（即作为线程）总是与启动异步操作的CUDA线程关联。异步操作使用同步对象来同步操作的完成。这样的同步对象可以由用户显式管理（例如，cuda::memcpy_async），也可以在库中隐式管理（例如，cooperative_groups::memcpy_async）。
同步对象可以是cuda::barrier或cuda::pipeline。这些对象在“异步屏障章节”和“使用cuda::pipeline进行异步数据复制章节”中有详细的解释。这些同步对象可以在不同的线程范围内使用。范围定义了可能使用同步对象与异步操作同步的线程集。下表定义了CUDA C++中可用的线程范围，以及可以与每个范围同步的线程。
| Thread Scope | Description |
|---|---|
| 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++的扩展来实现。
5.6 计算能力 Compute Capability
设备的计算能力用一个版本号表示，有时也被称为其“SM版本”。这个版本号标识了GPU硬件支持的特性，应用程序在运行时使用它来确定当前GPU上可用的硬件特性和/或指令。
计算能力由一个主要修订号X和一个次要修订号Y组成，表示为X.Y。
具有相同主修订号的设备具有相同的核心架构。主修订号为9的设备是基于NVIDIA Hopper GPU架构的，为8的设备是基于NVIDIA Ampere GPU架构的，为7的设备是基于Volta架构的，为6的设备是基于Pascal架构的，为5的设备是基于Maxwell架构的，为3的设备是基于Kepler架构的。
次修订号对应于对核心架构的增量改进，可能包括新的特性。
Turing是计算能力为7.5的设备的架构，是基于Volta架构的增量更新。
CUDA启用的GPU列表包含所有启用CUDA的设备及其计算能力。每种计算能力的技术规格在计算能力中提供。Tesla和Fermi架构从CUDA 7.0和CUDA 9.0开始分别不再支持。
特定GPU的计算能力版本不应与CUDA版本（例如，CUDA 7.5、CUDA 8、CUDA 9）混淆，后者是CUDA软件平台的版本。CUDA平台被应用开发者用来创建可以在许多代的GPU架构上运行的应用，包括尚未发明的未来GPU架构。虽然新版本的CUDA平台通常通过支持该架构的计算能力版本来增加对新GPU架构的本地支持，但新版本的CUDA平台通常也包括独立于硬件生成的软件特性。
