N.2.2.1. GPU Exclusive Access To Managed Memory 为确保 6.x 之前的 GPU 架构的一致性，统一内存编程模型限制 CPU 和 GPU 同时执行时的数据访问。GPU 对所有托管数据具有独占访问权，即使特定内核未使用数据。`cudaMemcpy*()` 或 `cudaMemset*()` 与托管数据一起使用时，系统可能选择主机或设备...
CUDA 编程手册系列 附录N – CUDA的统一内存五
### N.2.2. Coherency and Concurrency
在计算能力低于 6.x 的设备上，同时访问托管内存是不可行的。当 CPU 在 GPU 内核处于活动状态时访问统一内存分配，不能保证一致性。然而，计算能力支持至 6.x 的设备允许 CPU 和 GPU 通过新的页面错误机制同时访问统一内存分配。程序可通过检查 `concurrentManagedAccess` 属性来确定设备是否支持托管内存的并发访问。开发人员需确保正确同步，以防处理器间的数据冲突。
#### N.2.2.1. GPU Exclusive Access To Managed Memory
为确保 6.x 之前的 GPU 架构的一致性，统一内存编程模型限制 CPU 和 GPU 同时执行时的数据访问。GPU 对所有托管数据具有独占访问权，即使特定内核未使用数据。`cudaMemcpy*()` 或 `cudaMemset*()` 与托管数据一起使用时，系统可能选择主机或设备访问源或目标，限制并发 CPU 访问该数据。并发 CPU 访问托管数据在 6.x 之前的架构上无效。程序必须在访问托管数据前显式同步。
#### N.2.2.2. Explicit Synchronization and Logical GPU Activity
即使内核快速运行并在 CPU 接触托管数据前完成，仍需显式同步。统一内存使用逻辑活动来确定 GPU 是否空闲。在 GPU 处于活动状态时，CPU 访问托管数据需同步。任何保证 GPU 工作完成的函数调用都有效，包括 `cudaDeviceSynchronize()` 和 `cudaStreamSynchronize()` 等。在流或事件上同步可推断其他流的完成。CPU 从流回调访问托管数据需确保 GPU 上无其他可能访问的数据流处于活动状态。
有几个重要的注意事项：
- 在 GPU 处于活动状态时，始终允许 CPU 访问非托管零拷贝数据。
- GPU 在运行任何内核时被认为是活动的，即使该内核不使用托管数据。如果内核可能使用数据，则禁止访问，除非设备属性 `concurrentManagedAccess` 为 1。
- 托管内存的并发 GPU 间访问没有任何限制，除了适用于非托管内存的多 GPU 访问外。
- 并发 GPU 内核访问托管数据没有任何限制，允许 GPU 内核之间的竞争。
以上内容强调了在不同 GPU 架构下托管内存访问的限制与同步要求。通过正确同步和检查设备属性，开发者可以有效利用并发访问托管内存的潜力。2024-10-23
