4.14. L2 缓存控制
当 CUDA 核函数反复访问全局内存中的某个数据区域时，此类数据访问可视为持久访问。另一方面，如果数据只被访问一次，则此类数据访问可视为流式访问。
计算能力为 8.0 及以上的设备能够影响数据在 L2 缓存中的持久性，从而有可能以更高的带宽和更低的延迟访问全局内存。
此功能通过两个主要 API 提供：
CUDA 运行时 API（从 CUDA 11.0 开始）提供了对 L2 缓存持久性的编程控制。
libcu++库中的
cuda::annotated_ptrAPI（从 CUDA 11.5 开始）为 CUDA 核函数中的指针标注内存访问属性，以实现类似的效果..
以下各节重点介绍 CUDA 运行时 API。有关cuda::annotated_ptr方法的详细信息，请参阅libcu++ 文档。
4.14.1. 为持久访问预留 L2 缓存
可以预留一部分 L2 缓存，用于对全局内存的持久数据访问。持久访问优先使用这部分预留的 L2 缓存，而对全局内存的普通访问或流式访问只有在持久访问未使用这部分 L2 缓存时，才能使用它。
为持久访问预留的 L2 缓存大小可以在限制范围内调整：
cudaGetDeviceProperties(&prop, device_id);
size_t size = min(int(prop.l2CacheSize * 0.75), prop.persistingL2CacheMaxSize);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size); /* set-aside 3/4 of L2 cache for persisting accesses or the max allowed*/
当 GPU 配置为多实例 GPU（MIG）模式时，L2 缓存预留功能会被禁用。
使用多进程服务（MPS）时，不能通过cudaDeviceSetLimit更改 L2 缓存的预留大小。预留大小只能在 MPS 服务器启动时通过环境变量CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT指定。
4.14.2. 持久访问的 L2 策略
访问策略窗口指定全局内存中的一个连续区域，以及对该区域内的访问在 L2 缓存中的持久性属性。
下面的代码示例展示了如何使用 CUDA 流设置 L2 持久访问窗口。
CUDA 流示例
cudaStreamAttrValue stream_attribute; // Stream level attributes data structure
stream_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(ptr); // Global Memory data pointer
stream_attribute.accessPolicyWindow.num_bytes = num_bytes; // Number of bytes for persistence access.
// (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)
stream_attribute.accessPolicyWindow.hitRatio = 0.6; // Hint for cache hit ratio
stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting; // Type of access property on cache hit
stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; // Type of access property on cache miss.
//Set the attributes to a CUDA stream of type cudaStream_t
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
当核函数随后在 CUDAstream中执行时，对全局内存范围[ptr..ptr+num_bytes)内的内存访问比对其他全局内存位置的访问更有可能持久保留在 L2 缓存中。
也可以为 CUDA 图核函数节点设置 L2 持久性，如下例所示：
CUDA GraphKernelNode 示例
cudaKernelNodeAttrValue node_attribute; // Kernel level attributes data structure
node_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(ptr); // Global Memory data pointer
node_attribute.accessPolicyWindow.num_bytes = num_bytes; // Number of bytes for persistence access.
// (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)
node_attribute.accessPolicyWindow.hitRatio = 0.6; // Hint for cache hit ratio
node_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting; // Type of access property on cache hit
node_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; // Type of access property on cache miss.
//Set the attributes to a CUDA Graph Kernel node of type cudaGraphNode_t
cudaGraphKernelNodeSetAttribute(node, cudaKernelNodeAttributeAccessPolicyWindow, &node_attribute);
可以使用hitRatio参数指定获得hitProp属性的访问所占的比例。在上述两个示例中，全局内存区域[ptr..ptr+num_bytes)中 60% 的内存访问具有持久属性，40% 的内存访问具有流式属性。哪些具体的内存访问被归类为持久访问（hitProp）是随机的，概率约为hitRatio；概率分布取决于硬件架构和内存范围。
例如，如果 L2 预留缓存大小为 16KB，且accessPolicyWindow中的num_bytes为 32KB：
当
hitRatio为 0.5 时，硬件会从 32KB 窗口中随机选取 16KB，将其指定为持久数据，并缓存在预留的 L2 缓存区域中。当
hitRatio为 1.0 时，硬件会尝试将整个 32KB 窗口缓存在预留的 L2 缓存区域中。由于预留区域小于窗口，缓存行将被逐出，以便将 32KB 数据中最近使用的 16KB 保留在 L2 缓存的预留部分中。
因此，可以使用hitRatio来避免缓存行抖动，并从总体上减少移入和移出 L2 缓存的数据量。
可以使用小于 1.0 的hitRatio值，手动控制并发 CUDA 流的不同accessPolicyWindow能够在 L2 中缓存的数据量。例如，假设 L2 预留缓存大小为 16KB；两个不同 CUDA 流中的两个并发核函数各有一个 16KB 的accessPolicyWindow，且hitRatio值均为 1.0，它们在竞争共享的 L2 资源时，可能会逐出彼此的缓存行。但是，如果这两个accessPolicyWindows的 hitRatio 值均为 0.5，那么它们逐出自身或彼此的持久缓存行的可能性就会降低。
4.14.3. L2 访问属性
针对不同的全局内存数据访问，定义了三种访问属性：
cudaAccessPropertyStreaming：具有流式属性的内存访问不太可能持久保留在 L2 缓存中，因为这些访问会被优先逐出。cudaAccessPropertyPersisting：具有持久属性的内存访问更有可能持久保留在 L2 缓存中，因为这些访问会被优先保留在 L2 缓存的预留部分中。cudaAccessPropertyNormal：此访问属性会将先前应用的持久访问属性强制重置为普通状态。先前 CUDA 核函数中具有持久属性的内存访问，可能在其预期用途结束后很长时间内仍保留在 L2 缓存中。这种使用后仍持久保留的行为，会减少后续不使用持久属性的核函数可用的 L2 缓存量。使用cudaAccessPropertyNormal属性重置访问属性窗口，会移除先前访问的持久（优先保留）状态，就像先前的访问未带有访问属性一样。
4.14.4. L2 持久性示例
以下示例展示了如何为持久访问预留 L2 缓存，如何在 CUDA 核函数中通过 CUDA 流使用预留的 L2 缓存，以及如何随后重置 L2 缓存。
cudaStream_t stream;
cudaStreamCreate(&stream); // Create CUDA stream
cudaDeviceProp prop; // CUDA device properties variable
cudaGetDeviceProperties( &prop, device_id); // Query GPU properties
size_t size = min( int(prop.l2CacheSize * 0.75) , prop.persistingL2CacheMaxSize );
cudaDeviceSetLimit( cudaLimitPersistingL2CacheSize, size); // set-aside 3/4 of L2 cache for persisting accesses or the max allowed
size_t window_size = min(prop.accessPolicyMaxWindowSize, num_bytes); // Select minimum of user defined num_bytes and max window size.
cudaStreamAttrValue stream_attribute; // Stream level attributes data structure
stream_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(data1); // Global Memory data pointer
stream_attribute.accessPolicyWindow.num_bytes = window_size; // Number of bytes for persistence access
stream_attribute.accessPolicyWindow.hitRatio = 0.6; // Hint for cache hit ratio
stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting; // Persistence Property
stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; // Type of access property on cache miss
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute); // Set the attributes to a CUDA Stream
for(int i = 0; i < 10; i++) {
cuda_kernelA<<<grid_size,block_size,0,stream>>>(data1); // This data1 is used by a kernel multiple times
} // [data1 + num_bytes) benefits from L2 persistence
cuda_kernelB<<<grid_size,block_size,0,stream>>>(data1); // A different kernel in the same stream can also benefit
// from the persistence of data1
stream_attribute.accessPolicyWindow.num_bytes = 0; // Setting the window size to 0 disable it
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute); // Overwrite the access policy attribute to a CUDA Stream
cudaCtxResetPersistingL2Cache(); // Remove any persistent lines in L2
cuda_kernelC<<<grid_size,block_size,0,stream>>>(data2); // data2 can now benefit from full L2 in normal mode
4.14.5. 将 L2 访问重置为普通状态
来自先前 CUDA 核函数的持久 L2 缓存行，可能在使用后很长时间内仍持久保留在 L2 中。因此，要让流式或普通内存访问以普通优先级使用 L2 缓存，将 L2 缓存重置为普通状态很重要。有三种方法可以将持久访问重置为普通状态。
使用访问属性
cudaAccessPropertyNormal重置先前的持久内存区域。通过调用
cudaCtxResetPersistingL2Cache()，将所有持久 L2 缓存行重置为普通状态。未被访问的缓存行最终会自动重置为普通状态。强烈不建议依赖自动重置，因为自动重置发生所需的时间长度是不确定的。
4.14.6. 管理 L2 预留缓存的利用情况
在不同 CUDA 流中并发执行的多个 CUDA 核函数，其所在的流可能各自分配了不同的访问策略窗口。但是，L2 预留缓存部分由所有这些并发 CUDA 核函数共享。因此，这部分预留缓存的净使用量是所有并发核函数各自使用量的总和。当持久访问的数据量超过 L2 预留缓存容量时，将内存访问指定为持久访问所带来的收益就会降低。
要管理 L2 缓存预留部分的利用情况，应用程序必须考虑以下因素：
L2 预留缓存的大小。
可能并发执行的 CUDA 核函数。
所有可能并发执行的 CUDA 核函数的访问策略窗口。
何时以及如何重置 L2，才能使普通访问或流式访问以相同的优先级使用先前预留的 L2 缓存。
4.14.7. 查询 L2 缓存属性
与 L2 缓存相关的属性是cudaDeviceProp结构体的一部分，可以使用 CUDA 运行时 APIcudaGetDeviceProperties进行查询
CUDA 设备属性包括：
l2CacheSize：GPU 上可用的 L2 缓存量。persistingL2CacheMaxSize：可为持久内存访问预留的最大 L2 缓存量。accessPolicyMaxWindowSize：访问策略窗口的最大大小。
4.14.8. 控制为持久内存访问预留的 L2 缓存大小
为持久内存访问预留的 L2 缓存大小可使用 CUDA 运行时 APIcudaDeviceGetLimit查询，并使用 CUDA 运行时 APIcudaDeviceSetLimit将其作为cudaLimit进行设置。设置此限制时的最大值为cudaDeviceProp::persistingL2CacheMaxSize。
enum cudaLimit {
/* other fields not shown */
cudaLimitPersistingL2CacheSize
};