CUDA 流属性与 L2 访问策略窗口实战深入解析 cuda-samples simpleAttributes 示例【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文以 CUDA Samples 仓库中 cpp/0_Introduction/simpleAttributes 示例为核心系统讲解 CUDA Runtime API 中**流属性Stream Attributes的用法重点聚焦影响 L2 缓存局部性L2 locality的访问策略窗口Access Policy Window**机制。读完本文你将掌握cudaStreamSetAttribute与cudaAccessPolicyWindow的完整配置流程、持久化 L2 缓存Persisting L2 Cache的启用方法以及该特性在 AmpereCompute Capability 8.0及以上架构上的性能收益原理与实测验证手段。示例概述与核心价值simpleAttributes是一个极简的 CUDA Runtime API 入门示例其唯一目的就是演示如何通过流属性控制数据在 L2 缓存中的驻留策略。根据 simpleAttributes/README.md 的官方描述This CUDA Runtime API sample is a very basic example that implements how to use the stream attributes that affect L2 locality. Performance improvement due to use of L2 access policy window can only be noticed on Compute capability 8.0 or higher.关键信息有两点功能定位演示影响 L2 局部性的流属性的使用方法属于入门级Introduction示例。硬件前提L2 访问策略窗口带来的性能提升仅在 Compute Capability 8.0 及以上即 Ampere、Hopper、Ada Lovelace、Blackwell 架构才能观察到。这是因为持久化 L2 缓存Persisting L2与访问策略窗口是这些新架构才引入的硬件特性。该示例的核心思想是同一份数据在不设置访问策略时由普通 L2 缓存行处理可能被频繁驱逐而设置持久化访问策略后可以钉在 L2 中从而减少 DRAM 带宽消耗、提升重复访问命中率。支持的平台与前置条件支持的硬件架构SM Architectures根据 README 的 Supported SM Architectures 一节该示例支持从 Maxwell 到 Blackwell 的全系列主流 SM 架构SM 5.0、SM 5.2、SM 5.3、SM 6.0、SM 6.1、SM 7.0、SM 7.2、SM 7.5、SM 8.0、SM 8.6、SM 8.7、SM 8.9、SM 9.0需要注意的是编译层面支持这些架构但性能收益层面持久化 L2 窗口生效只有 SM 8.0 及以上的设备才能体现。源码中对此有显式的运行期校验见下文设备能力检测小节。支持的操作系统与 CPU 架构操作系统Linux、WindowsCPU 架构x86_64、armv7l前置条件与仓库内所有 CUDA Samples 一致唯一硬性前置条件是Download and install the CUDA Toolkit for your corresponding platform.即先安装与平台匹配的 CUDA Toolkit本仓库对应 CUDA Toolkit 13.3见 根目录 README.md 的版本说明然后即可编译运行该示例。涉及的 CUDA Runtime API 清单README 明确列出了本示例用到的全部 CUDA Runtime API这些 API 在源码 simpleAttributes.cu 中均有实际调用可逐一对照API在本示例中的用途cudaGetDeviceProperties查询设备属性特别是l2CacheSize与persistingL2CacheMaxSizecudaStreamCreate创建将要挂载访问策略窗口的 CUDA 流cudaDeviceSetLimit将持久化 L2 缓存上限设置为设备支持的最大值cudaLimitPersistingL2CacheSizecudaStreamSetAttribute核心 API将访问策略窗口cudaStreamAttributeAccessPolicyWindow绑定到流cudaMallocHost/cudaFreeHost分配/释放页锁定主机内存用于初始化数据cudaMalloc/cudaFree分配/释放设备全局内存数据缓冲与垃圾缓冲cudaMemcpyAsync在流上异步将主机数据拷贝到设备cudaStreamSynchronize同步流等待拷贝与内核执行完成cudaCtxResetPersistingL2Cache重置当前上下文中已有的持久化 L2 缓存行源码中实际使用说明README 列出的 API 清单中未包含cudaCtxResetPersistingL2Cache但该 API 在源码 simpleAttributes.cu 中确有调用用于在设置新窗口前清空上一次遗留的持久化行保证测试的公平性。源码逐段解析访问策略窗口的完整生命周期示例主体位于 simpleAttributes.cu共约 208 行。下面按执行流程拆解。4.1 访问策略窗口的默认初始化源码开篇定义了一个辅助函数initAccessPolicyWindow返回一个全默认的窗口结构体cudaAccessPolicyWindow initAccessPolicyWindow(void) { cudaAccessPolicyWindow accessPolicyWindow {0}; accessPolicyWindow.base_ptr (void *)0; accessPolicyWindow.num_bytes 0; accessPolicyWindow.hitRatio 0.f; accessPolicyWindow.hitProp cudaAccessPropertyNormal; accessPolicyWindow.missProp cudaAccessPropertyStreaming; return accessPolicyWindow; }cudaAccessPolicyWindow是 CUDA Runtime API 中描述 L2 访问策略的核心结构体五个字段含义如下字段类型含义base_ptrvoid*策略窗口覆盖的全局内存起始地址num_bytessize_t窗口覆盖的字节数0 表示不生效hitRatiofloat命中比例取值范围 [0.0, 1.0]表示窗口内有多少比例的数据访问按hitProp处理hitPropcudaAccessProperty命中时采用的访问属性Normal / Streaming / PersistingmissPropcudaAccessProperty未命中窗口比例之外的访问采用的访问属性这里体现了一个容易被忽略的设计细节默认窗口把missProp设为cudaAccessPropertyStreaming流式属性即默认将窗口内未命中的访问标记为流式。cudaAccessProperty枚举的三个取值语义为cudaAccessPropertyNormal普通访问遵循常规 L2 替换策略cudaAccessPropertyStreaming流式访问数据倾向被标记为少复用减少对 L2 的污染cudaAccessPropertyPersisting持久化访问数据行可驻留于持久化 L2 分区供后续重复访问命中。4.2 设备选择与能力检测runTest中首先通过findCudaDevice(argc, argv)来自 Common/helper_cuda.h 的辅助函数选择命令行指定或计算能力最高的设备然后查询设备属性int devID findCudaDevice(argc, (const char **)argv); checkCudaErrors(cudaGetDeviceProperties(deviceProp, devID)); dim3 blocks(deviceProp.maxGridSize[1], 1); // Make sure device the l2 optimization if (deviceProp.persistingL2CacheMaxSize 0) { printf(Waiving execution as device %d does not support persisting L2 Caching\n, devID); exit(EXIT_WAIVED); }这段代码是运行期硬件能力门槛cudaDeviceProp.persistingL2CacheMaxSize表示该设备持久化 L2 缓存的最大字节数。如果为 0说明设备不支持持久化 L2 缓存即 Compute Capability 低于 8.0程序直接以EXIT_WAIVED值 2定义于 Common/helper_cuda.h退出并提示放弃执行。这与 README 中性能提升仅限 CC 8.0的说明完全对应。4.3 创建流并扩大持久化 L2 配额checkCudaErrors(cudaStreamCreate(stream)); // Set the amount of l2 cache that will be persisting to maximum the device // can support checkCudaErrors(cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, deviceProp.persistingL2CacheMaxSize));cudaStreamCreate创建一条普通流后续所有操作数据拷贝、内核启动都在该流上执行以便访问策略窗口随流生效cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, ...)把持久化 L2 缓存配额设置为设备支持的最大值。注意持久化 L2 是设备级资源配额必须先分配额度流上的窗口策略才能实际将行驻留其中。4.4 缓冲区规划如何构造窗口命中/未命中对照实验示例精妙之处在于用两块缓冲区构造了对照实验bigDataSize (deviceProp.l2CacheSize * 4) / sizeof(int); dataSize (deviceProp.l2CacheSize / 4) / sizeof(int); checkCudaErrors(cudaMallocHost(dataHostPointer, dataSize * sizeof(int))); checkCudaErrors(cudaMallocHost(bigDataHostPointer, bigDataSize * sizeof(int)));data缓冲小大小仅为l2CacheSize / 4个 int即设备 L2 总容量的 1/4后续将作为访问策略窗口的覆盖对象bigData缓冲大源码中变量名trash大小是l2CacheSize * 4个 int即 L2 容量的 4 倍远超 L2 容量用于制造必然挤占 L2的访问压力。随后用cudaMallocHost分配页锁定主机内存完成初始化把data初始化为递增序列、bigData初始化为倒序序列再通过cudaMemcpyAsync在流上异步拷入设备全局内存。使用页锁定内存 异步拷贝的组合可以避免阻塞流保证策略窗口从一开始就作用于该流上的全部内存操作。4.5 配置并挂载访问策略窗口核心// Make a window for the buffer of interest accessPolicyWindow.base_ptr (void *)dataDevicePointer; accessPolicyWindow.num_bytes dataSize * sizeof(int); accessPolicyWindow.hitRatio 1.f; accessPolicyWindow.hitProp cudaAccessPropertyPersisting; accessPolicyWindow.missProp cudaAccessPropertyNormal; streamAttrValue.accessPolicyWindow accessPolicyWindow; // Assign window to stream checkCudaErrors(cudaStreamSetAttribute(stream, streamAttrID, streamAttrValue));这里完成了三个关键设定窗口范围base_ptr dataDevicePointer、num_bytes dataSize * sizeof(int)即窗口精确覆盖小而重要的data缓冲区命中策略hitRatio 1.f、hitProp cudaAccessPropertyPersisting即窗口内 100% 的数据访问都按持久化处理使这些行驻留在持久化 L2 分区未命中策略missProp cudaAccessPropertyNormal窗口之外或超出比例的访问走普通 L2 路径。streamAttrID被赋值为cudaStreamAttributeAccessPolicyWindow这是cudaStreamAttrID枚举中用于标识访问策略窗口属性的取值streamAttrValue则是cudaStreamAttrValue联合体通过其accessPolicyWindow成员携带窗口结构体。挂载之前示例还调用了一次// Demote any previous persisting lines checkCudaErrors(cudaCtxResetPersistingL2Cache());cudaCtxResetPersistingL2Cache()用于降级demote当前上下文中此前遗留的所有持久化 L2 行确保新窗口生效前 L2 处于干净状态避免历史行影响测试结果。4.6 内核设计制造命中与未命中的对称访问static __global__ void kernCacheSegmentTest(int *data, int dataSize, int *trash, int bigDataSize, int hitCount) { __shared__ unsigned int hit; int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; int tID row * blockDim.y col; uint32_t psRand tID; atomicExch(hit, 0); __syncthreads(); while (hit hitCount) { psRand ^ psRand 13; psRand ^ psRand 17; psRand ^ psRand 5; int idx tID - psRand; if (idx 0) { idx -idx; } if ((tID % 2) 0) { data[psRand % dataSize] data[psRand % dataSize] data[idx % dataSize]; } else { trash[psRand % bigDataSize] trash[psRand % bigDataSize] trash[idx % bigDataSize]; } atomicAdd(hit, 1); } }内核的对照逻辑非常清晰伪随机访问每轮循环用 xorshift 三段式13、17、5更新伪随机种子psRand模拟看似随机、实则重复的访存模式偶数线程tID % 2 0反复读写窗口内的data缓冲约 L2 的 1/4容量足够装进持久化分区这些访问落在访问策略窗口内应被持久化驻留奇数线程tID % 2 ! 0反复读写远超 L2 容量的trash缓冲4 倍 L2 大小制造大量的 L2 未命中与 DRAM 流量持续冲刷L2共享内存计数用atomicExch/atomicAdd在 block 内协同控制循环次数hitCount 0xAFFFF约 72 万次保证所有线程执行大致相同的访存工作量。通过一半线程访问窗口内数据、一半线程访问窗口外大缓冲的设计示例可以在同一内核中同时观察持久化窗口内数据的命中情况以及普通访问被大缓冲挤占的对比效果。这也是为什么该示例能直观体现访问策略窗口价值的根本原因。4.7 启动、计时与清理checkCudaErrors(cudaStreamSynchronize(stream)); kernCacheSegmentTestblocks, threads, 0, stream( dataDevicePointer, dataSize, bigDataDevicePointer, bigDataSize, 0xAFFFF); checkCudaErrors(cudaStreamSynchronize(stream)); getLastCudaError(Kernel execution failed);网格配置threads(32, 32)每 block 1024 线程blocks(deviceProp.maxGridSize[1], 1)按设备最大网格维度生成内核在挂载了访问策略窗口的stream上启动blocks, threads, 0, stream因此窗口策略对内核访存全程生效getLastCudaError来自 Common/helper_cuda.h检查内核执行错误程序用sdkCreateTimer/sdkStartTimer/sdkStopTimer/sdkGetTimerValue来自 Common/helper_functions.h统计整体处理耗时并打印Processing time: ... (ms)随后释放主机与设备内存、退出。从整个流程可以看到README 中列出的全部 API 都被串联在一条清晰的主线上查设备 → 校验能力 → 建流 → 扩配额 → 配窗口 → 挂流 → 跑内核 → 计时清理。CMake 构建配置解析该示例的构建文件 cpp/0_Introduction/simpleAttributes/CMakeLists.txt 结构简洁值得注意的要点cmake_minimum_required(VERSION 3.20) project(simpleAttributes LANGUAGES C CXX CUDA) find_package(CUDAToolkit REQUIRED) set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120) include_directories(../../../Common) add_executable(simpleAttributes simpleAttributes.cu) set_target_properties(simpleAttributes PROPERTIES CUDA_SEPARABLE_COMPILATION ON) include(${CMAKE_CURRENT_SOURCE_DIR}/../../../cmake/InstallSamples.cmake) setup_samples_install()CMake 最低版本 3.20与仓库根 CMakeLists.txt 的要求一致CMAKE_CUDA_ARCHITECTURES覆盖 SM 7.5 至 SM 12.0与 README 的支持矩阵对应默认不包含 5.x 旧架构头文件搜索路径../../../Common指向仓库的 Common 公共目录示例用到的helper_cuda.h、helper_functions.h均在此处CUDA_SEPARABLE_COMPILATION ON启用可分离编译setup_samples_install()来自 cmake/InstallSamples.cmake负责安装配置。该示例通过根目录 CMakeLists.txt → cpp/CMakeLists.txt → cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(simpleAttributes)被纳入整体构建。按照 根目录 README.md 的说明构建方式为# 在仓库根目录 mkdir build cd build cmake .. make -j$(nproc)也可直接在simpleAttributes子目录内单独配置构建。Windows 下则可通过 Visual Studio 直接导入或使用cmake .. -G Visual Studio 16 2019 -A x64生成解决方案。运行与结果解读运行编译出的simpleAttributes可执行文件Linux 下位于构建目录程序将自动选择合适设备并执行若设备不支持持久化 L2如 Maxwell/Pascal 等 CC 8.0 的 GPU会输出Waiving execution as device id does not support persisting L2 Caching并以EXIT_WAIVED退出这并非错误而是硬件能力不足的友好提示若设备支持CC 8.0将打印Starting...、内核执行信息以及Processing time: t (ms)的处理耗时。通过命令行参数可指定设备例如./simpleAttributes -device0设备选择逻辑由findCudaDeviceCommon/helper_cuda.h实现。由于该示例在单次运行中同时构造了窗口命中与窗口外冲刷两类访问读者可以在此基础上对比将accessPolicyWindow.hitProp从cudaAccessPropertyPersisting改为cudaAccessPropertyStreaming或直接将num_bytes置 0 关闭窗口观察耗时变化即可直观量化访问策略窗口的收益。关键结论与使用建议访问策略窗口是流级属性通过cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, value)挂载作用于该流上的全部内存访问适合与cudaMemcpyAsync 内核启动组成完整的持久化数据流水线配额先行使用持久化属性前必须用cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, ...)分配设备级持久化 L2 配额cudaDeviceProp.persistingL2CacheMaxSize给出上限运行期校验不可省示例在运行期检查persistingL2CacheMaxSize 0来规避不支持设备生产代码同样应保留该检查并注意 README 明确指出的前提——性能提升仅在 Compute Capability 8.0 及以上可观察到窗口命中比例可调hitRatio介于 0.01.0hitProp/missProp组合Normal / Streaming / Persisting决定了窗口内外的 L2 行处置策略是调优数据驻留的核心旋钮测试公平性cudaCtxResetPersistingL2Cache()在设置新窗口前降级旧持久化行这是做 A/B 性能对比时容易被遗漏但至关重要的细节。该示例定位是入门级0_Introduction代码量小但完整覆盖了从能力检测、配额分配、窗口配置到内核验证的全链路是学习 CUDA L2 缓存精细控制的最佳起点如需进一步了解相关概念可继续阅读仓库 cpp/0_Introduction 下的其他流相关示例如 simpleStreams、simpleMultiCopy进行横向对照。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CUDA 流属性与 L2 访问策略窗口实战：深入解析 cuda-samples simpleAttributes 示例
CUDA 流属性与 L2 访问策略窗口实战深入解析 cuda-samples simpleAttributes 示例【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples导读本文以 CUDA Samples 仓库中 cpp/0_Introduction/simpleAttributes 示例为核心系统讲解 CUDA Runtime API 中**流属性Stream Attributes的用法重点聚焦影响 L2 缓存局部性L2 locality的访问策略窗口Access Policy Window**机制。读完本文你将掌握cudaStreamSetAttribute与cudaAccessPolicyWindow的完整配置流程、持久化 L2 缓存Persisting L2 Cache的启用方法以及该特性在 AmpereCompute Capability 8.0及以上架构上的性能收益原理与实测验证手段。示例概述与核心价值simpleAttributes是一个极简的 CUDA Runtime API 入门示例其唯一目的就是演示如何通过流属性控制数据在 L2 缓存中的驻留策略。根据 simpleAttributes/README.md 的官方描述This CUDA Runtime API sample is a very basic example that implements how to use the stream attributes that affect L2 locality. Performance improvement due to use of L2 access policy window can only be noticed on Compute capability 8.0 or higher.关键信息有两点功能定位演示影响 L2 局部性的流属性的使用方法属于入门级Introduction示例。硬件前提L2 访问策略窗口带来的性能提升仅在 Compute Capability 8.0 及以上即 Ampere、Hopper、Ada Lovelace、Blackwell 架构才能观察到。这是因为持久化 L2 缓存Persisting L2与访问策略窗口是这些新架构才引入的硬件特性。该示例的核心思想是同一份数据在不设置访问策略时由普通 L2 缓存行处理可能被频繁驱逐而设置持久化访问策略后可以钉在 L2 中从而减少 DRAM 带宽消耗、提升重复访问命中率。支持的平台与前置条件支持的硬件架构SM Architectures根据 README 的 Supported SM Architectures 一节该示例支持从 Maxwell 到 Blackwell 的全系列主流 SM 架构SM 5.0、SM 5.2、SM 5.3、SM 6.0、SM 6.1、SM 7.0、SM 7.2、SM 7.5、SM 8.0、SM 8.6、SM 8.7、SM 8.9、SM 9.0需要注意的是编译层面支持这些架构但性能收益层面持久化 L2 窗口生效只有 SM 8.0 及以上的设备才能体现。源码中对此有显式的运行期校验见下文设备能力检测小节。支持的操作系统与 CPU 架构操作系统Linux、WindowsCPU 架构x86_64、armv7l前置条件与仓库内所有 CUDA Samples 一致唯一硬性前置条件是Download and install the CUDA Toolkit for your corresponding platform.即先安装与平台匹配的 CUDA Toolkit本仓库对应 CUDA Toolkit 13.3见 根目录 README.md 的版本说明然后即可编译运行该示例。涉及的 CUDA Runtime API 清单README 明确列出了本示例用到的全部 CUDA Runtime API这些 API 在源码 simpleAttributes.cu 中均有实际调用可逐一对照API在本示例中的用途cudaGetDeviceProperties查询设备属性特别是l2CacheSize与persistingL2CacheMaxSizecudaStreamCreate创建将要挂载访问策略窗口的 CUDA 流cudaDeviceSetLimit将持久化 L2 缓存上限设置为设备支持的最大值cudaLimitPersistingL2CacheSizecudaStreamSetAttribute核心 API将访问策略窗口cudaStreamAttributeAccessPolicyWindow绑定到流cudaMallocHost/cudaFreeHost分配/释放页锁定主机内存用于初始化数据cudaMalloc/cudaFree分配/释放设备全局内存数据缓冲与垃圾缓冲cudaMemcpyAsync在流上异步将主机数据拷贝到设备cudaStreamSynchronize同步流等待拷贝与内核执行完成cudaCtxResetPersistingL2Cache重置当前上下文中已有的持久化 L2 缓存行源码中实际使用说明README 列出的 API 清单中未包含cudaCtxResetPersistingL2Cache但该 API 在源码 simpleAttributes.cu 中确有调用用于在设置新窗口前清空上一次遗留的持久化行保证测试的公平性。源码逐段解析访问策略窗口的完整生命周期示例主体位于 simpleAttributes.cu共约 208 行。下面按执行流程拆解。4.1 访问策略窗口的默认初始化源码开篇定义了一个辅助函数initAccessPolicyWindow返回一个全默认的窗口结构体cudaAccessPolicyWindow initAccessPolicyWindow(void) { cudaAccessPolicyWindow accessPolicyWindow {0}; accessPolicyWindow.base_ptr (void *)0; accessPolicyWindow.num_bytes 0; accessPolicyWindow.hitRatio 0.f; accessPolicyWindow.hitProp cudaAccessPropertyNormal; accessPolicyWindow.missProp cudaAccessPropertyStreaming; return accessPolicyWindow; }cudaAccessPolicyWindow是 CUDA Runtime API 中描述 L2 访问策略的核心结构体五个字段含义如下字段类型含义base_ptrvoid*策略窗口覆盖的全局内存起始地址num_bytessize_t窗口覆盖的字节数0 表示不生效hitRatiofloat命中比例取值范围 [0.0, 1.0]表示窗口内有多少比例的数据访问按hitProp处理hitPropcudaAccessProperty命中时采用的访问属性Normal / Streaming / PersistingmissPropcudaAccessProperty未命中窗口比例之外的访问采用的访问属性这里体现了一个容易被忽略的设计细节默认窗口把missProp设为cudaAccessPropertyStreaming流式属性即默认将窗口内未命中的访问标记为流式。cudaAccessProperty枚举的三个取值语义为cudaAccessPropertyNormal普通访问遵循常规 L2 替换策略cudaAccessPropertyStreaming流式访问数据倾向被标记为少复用减少对 L2 的污染cudaAccessPropertyPersisting持久化访问数据行可驻留于持久化 L2 分区供后续重复访问命中。4.2 设备选择与能力检测runTest中首先通过findCudaDevice(argc, argv)来自 Common/helper_cuda.h 的辅助函数选择命令行指定或计算能力最高的设备然后查询设备属性int devID findCudaDevice(argc, (const char **)argv); checkCudaErrors(cudaGetDeviceProperties(deviceProp, devID)); dim3 blocks(deviceProp.maxGridSize[1], 1); // Make sure device the l2 optimization if (deviceProp.persistingL2CacheMaxSize 0) { printf(Waiving execution as device %d does not support persisting L2 Caching\n, devID); exit(EXIT_WAIVED); }这段代码是运行期硬件能力门槛cudaDeviceProp.persistingL2CacheMaxSize表示该设备持久化 L2 缓存的最大字节数。如果为 0说明设备不支持持久化 L2 缓存即 Compute Capability 低于 8.0程序直接以EXIT_WAIVED值 2定义于 Common/helper_cuda.h退出并提示放弃执行。这与 README 中性能提升仅限 CC 8.0的说明完全对应。4.3 创建流并扩大持久化 L2 配额checkCudaErrors(cudaStreamCreate(stream)); // Set the amount of l2 cache that will be persisting to maximum the device // can support checkCudaErrors(cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, deviceProp.persistingL2CacheMaxSize));cudaStreamCreate创建一条普通流后续所有操作数据拷贝、内核启动都在该流上执行以便访问策略窗口随流生效cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, ...)把持久化 L2 缓存配额设置为设备支持的最大值。注意持久化 L2 是设备级资源配额必须先分配额度流上的窗口策略才能实际将行驻留其中。4.4 缓冲区规划如何构造窗口命中/未命中对照实验示例精妙之处在于用两块缓冲区构造了对照实验bigDataSize (deviceProp.l2CacheSize * 4) / sizeof(int); dataSize (deviceProp.l2CacheSize / 4) / sizeof(int); checkCudaErrors(cudaMallocHost(dataHostPointer, dataSize * sizeof(int))); checkCudaErrors(cudaMallocHost(bigDataHostPointer, bigDataSize * sizeof(int)));data缓冲小大小仅为l2CacheSize / 4个 int即设备 L2 总容量的 1/4后续将作为访问策略窗口的覆盖对象bigData缓冲大源码中变量名trash大小是l2CacheSize * 4个 int即 L2 容量的 4 倍远超 L2 容量用于制造必然挤占 L2的访问压力。随后用cudaMallocHost分配页锁定主机内存完成初始化把data初始化为递增序列、bigData初始化为倒序序列再通过cudaMemcpyAsync在流上异步拷入设备全局内存。使用页锁定内存 异步拷贝的组合可以避免阻塞流保证策略窗口从一开始就作用于该流上的全部内存操作。4.5 配置并挂载访问策略窗口核心// Make a window for the buffer of interest accessPolicyWindow.base_ptr (void *)dataDevicePointer; accessPolicyWindow.num_bytes dataSize * sizeof(int); accessPolicyWindow.hitRatio 1.f; accessPolicyWindow.hitProp cudaAccessPropertyPersisting; accessPolicyWindow.missProp cudaAccessPropertyNormal; streamAttrValue.accessPolicyWindow accessPolicyWindow; // Assign window to stream checkCudaErrors(cudaStreamSetAttribute(stream, streamAttrID, streamAttrValue));这里完成了三个关键设定窗口范围base_ptr dataDevicePointer、num_bytes dataSize * sizeof(int)即窗口精确覆盖小而重要的data缓冲区命中策略hitRatio 1.f、hitProp cudaAccessPropertyPersisting即窗口内 100% 的数据访问都按持久化处理使这些行驻留在持久化 L2 分区未命中策略missProp cudaAccessPropertyNormal窗口之外或超出比例的访问走普通 L2 路径。streamAttrID被赋值为cudaStreamAttributeAccessPolicyWindow这是cudaStreamAttrID枚举中用于标识访问策略窗口属性的取值streamAttrValue则是cudaStreamAttrValue联合体通过其accessPolicyWindow成员携带窗口结构体。挂载之前示例还调用了一次// Demote any previous persisting lines checkCudaErrors(cudaCtxResetPersistingL2Cache());cudaCtxResetPersistingL2Cache()用于降级demote当前上下文中此前遗留的所有持久化 L2 行确保新窗口生效前 L2 处于干净状态避免历史行影响测试结果。4.6 内核设计制造命中与未命中的对称访问static __global__ void kernCacheSegmentTest(int *data, int dataSize, int *trash, int bigDataSize, int hitCount) { __shared__ unsigned int hit; int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; int tID row * blockDim.y col; uint32_t psRand tID; atomicExch(hit, 0); __syncthreads(); while (hit hitCount) { psRand ^ psRand 13; psRand ^ psRand 17; psRand ^ psRand 5; int idx tID - psRand; if (idx 0) { idx -idx; } if ((tID % 2) 0) { data[psRand % dataSize] data[psRand % dataSize] data[idx % dataSize]; } else { trash[psRand % bigDataSize] trash[psRand % bigDataSize] trash[idx % bigDataSize]; } atomicAdd(hit, 1); } }内核的对照逻辑非常清晰伪随机访问每轮循环用 xorshift 三段式13、17、5更新伪随机种子psRand模拟看似随机、实则重复的访存模式偶数线程tID % 2 0反复读写窗口内的data缓冲约 L2 的 1/4容量足够装进持久化分区这些访问落在访问策略窗口内应被持久化驻留奇数线程tID % 2 ! 0反复读写远超 L2 容量的trash缓冲4 倍 L2 大小制造大量的 L2 未命中与 DRAM 流量持续冲刷L2共享内存计数用atomicExch/atomicAdd在 block 内协同控制循环次数hitCount 0xAFFFF约 72 万次保证所有线程执行大致相同的访存工作量。通过一半线程访问窗口内数据、一半线程访问窗口外大缓冲的设计示例可以在同一内核中同时观察持久化窗口内数据的命中情况以及普通访问被大缓冲挤占的对比效果。这也是为什么该示例能直观体现访问策略窗口价值的根本原因。4.7 启动、计时与清理checkCudaErrors(cudaStreamSynchronize(stream)); kernCacheSegmentTestblocks, threads, 0, stream( dataDevicePointer, dataSize, bigDataDevicePointer, bigDataSize, 0xAFFFF); checkCudaErrors(cudaStreamSynchronize(stream)); getLastCudaError(Kernel execution failed);网格配置threads(32, 32)每 block 1024 线程blocks(deviceProp.maxGridSize[1], 1)按设备最大网格维度生成内核在挂载了访问策略窗口的stream上启动blocks, threads, 0, stream因此窗口策略对内核访存全程生效getLastCudaError来自 Common/helper_cuda.h检查内核执行错误程序用sdkCreateTimer/sdkStartTimer/sdkStopTimer/sdkGetTimerValue来自 Common/helper_functions.h统计整体处理耗时并打印Processing time: ... (ms)随后释放主机与设备内存、退出。从整个流程可以看到README 中列出的全部 API 都被串联在一条清晰的主线上查设备 → 校验能力 → 建流 → 扩配额 → 配窗口 → 挂流 → 跑内核 → 计时清理。CMake 构建配置解析该示例的构建文件 cpp/0_Introduction/simpleAttributes/CMakeLists.txt 结构简洁值得注意的要点cmake_minimum_required(VERSION 3.20) project(simpleAttributes LANGUAGES C CXX CUDA) find_package(CUDAToolkit REQUIRED) set(CMAKE_CUDA_ARCHITECTURES 75 80 86 87 89 90 100 110 120) include_directories(../../../Common) add_executable(simpleAttributes simpleAttributes.cu) set_target_properties(simpleAttributes PROPERTIES CUDA_SEPARABLE_COMPILATION ON) include(${CMAKE_CURRENT_SOURCE_DIR}/../../../cmake/InstallSamples.cmake) setup_samples_install()CMake 最低版本 3.20与仓库根 CMakeLists.txt 的要求一致CMAKE_CUDA_ARCHITECTURES覆盖 SM 7.5 至 SM 12.0与 README 的支持矩阵对应默认不包含 5.x 旧架构头文件搜索路径../../../Common指向仓库的 Common 公共目录示例用到的helper_cuda.h、helper_functions.h均在此处CUDA_SEPARABLE_COMPILATION ON启用可分离编译setup_samples_install()来自 cmake/InstallSamples.cmake负责安装配置。该示例通过根目录 CMakeLists.txt → cpp/CMakeLists.txt → cpp/0_Introduction/CMakeLists.txt 中的add_subdirectory(simpleAttributes)被纳入整体构建。按照 根目录 README.md 的说明构建方式为# 在仓库根目录 mkdir build cd build cmake .. make -j$(nproc)也可直接在simpleAttributes子目录内单独配置构建。Windows 下则可通过 Visual Studio 直接导入或使用cmake .. -G Visual Studio 16 2019 -A x64生成解决方案。运行与结果解读运行编译出的simpleAttributes可执行文件Linux 下位于构建目录程序将自动选择合适设备并执行若设备不支持持久化 L2如 Maxwell/Pascal 等 CC 8.0 的 GPU会输出Waiving execution as device id does not support persisting L2 Caching并以EXIT_WAIVED退出这并非错误而是硬件能力不足的友好提示若设备支持CC 8.0将打印Starting...、内核执行信息以及Processing time: t (ms)的处理耗时。通过命令行参数可指定设备例如./simpleAttributes -device0设备选择逻辑由findCudaDeviceCommon/helper_cuda.h实现。由于该示例在单次运行中同时构造了窗口命中与窗口外冲刷两类访问读者可以在此基础上对比将accessPolicyWindow.hitProp从cudaAccessPropertyPersisting改为cudaAccessPropertyStreaming或直接将num_bytes置 0 关闭窗口观察耗时变化即可直观量化访问策略窗口的收益。关键结论与使用建议访问策略窗口是流级属性通过cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, value)挂载作用于该流上的全部内存访问适合与cudaMemcpyAsync 内核启动组成完整的持久化数据流水线配额先行使用持久化属性前必须用cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, ...)分配设备级持久化 L2 配额cudaDeviceProp.persistingL2CacheMaxSize给出上限运行期校验不可省示例在运行期检查persistingL2CacheMaxSize 0来规避不支持设备生产代码同样应保留该检查并注意 README 明确指出的前提——性能提升仅在 Compute Capability 8.0 及以上可观察到窗口命中比例可调hitRatio介于 0.01.0hitProp/missProp组合Normal / Streaming / Persisting决定了窗口内外的 L2 行处置策略是调优数据驻留的核心旋钮测试公平性cudaCtxResetPersistingL2Cache()在设置新窗口前降级旧持久化行这是做 A/B 性能对比时容易被遗漏但至关重要的细节。该示例定位是入门级0_Introduction代码量小但完整覆盖了从能力检测、配额分配、窗口配置到内核验证的全链路是学习 CUDA L2 缓存精细控制的最佳起点如需进一步了解相关概念可继续阅读仓库 cpp/0_Introduction 下的其他流相关示例如 simpleStreams、simpleMultiCopy进行横向对照。【免费下载链接】cuda-samplesSamples for CUDA Developers which demonstrates features in CUDA Toolkit项目地址: https://gitcode.com/GitHub_Trending/cu/cuda-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考