5.3 内存层次 Memory Hierarchy
CUDA 线程在执行过程中可能会访问多个内存空间的数据。每个线程都有自己的私有本地内存。
每个线程块都有一个对块内所有线程可见的共享内存，并且其生命周期与块相同。线程块集群中的线程块可以对彼此的共享内存执行读、写和原子操作。所有线程都可以访问同一块全局内存。
此外，还有两个只读内存空间可以被所有线程访问：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间都针对不同的内存使用进行了优化。纹理内存也提供了不同的寻址模式，以及针对某些特定数据格式的数据过滤。
全局内存 (global)、常量内存 (constant) 和纹理内存 (texture) 空间在相同应用程序的内核启动间都是持久的。
5.4 异构编程
CUDA 编程模型假设 CUDA 线程在一个物理独立的设备上运行，该设备作为运行 C++ 程序的主机的协处理器。例如，当内核在 GPU 上运行，而 C++ 程序的其余部分在 CPU 上运行时，就是这种情况。
CUDA 编程模型还假设主机和设备在 DRAM 中分别维护自己的独立内存空间，分别称为主机内存和设备内存。因此，一个程序通过调用 CUDA 运行时来管理内核可以看到的全局内存、常量内存和纹理内存空间。这包括设备内存的分配和释放，以及主机和设备内存之间的数据传输。
统一内存提供了管理内存，以连接主机和设备的内存空间。管理内存可以作为一个统一、连贯的内存映像，通过一个共享的地址空间，从系统中的所有 CPU 和 GPU 访问。这一能力使设备内存能被过度订阅，并且可以大大简化转换应用程序的任务，因为它消除了在主机和设备之间明确镜像数据的需要。
串行代码在主机上执行，而并行代码在设备上执行
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。
