CUDA设备内存原子操作是否支持使用L1缓存？（cc7.5环境）

不存在配置错误或代码逻辑问题，CC 7.5（Turing架构）下全局内存的原子操作确实不会经过L1缓存，所有原子请求会直接下发到L2处理，你的NSight观测结果和硬件设计完全一致。

• Turing架构的SM上L1缓存与共享内存虽然共用硬件单元，但全局原子操作的硬件通路从设计上就绕过L1，直接路由到L2对应的原子处理单元，这就是你观测到L1缓存命中率、Global Atomic ALU L1命中率均为0的根本原因。
• 你查到的提到原子操作可经过L1的官方文档，描述的是CC 8.0及以上Ampere、Hopper等新架构的特性：从Ampere开始SM内新增了支持L1处理原子的硬件路径，命中L1的原子操作无需访问L2即可完成，这个特性在CC7.5硬件上不存在。
• 早期开发者论坛提到的“设备内存原子始终通过L2完成”的结论，对CC7.5及更早的所有架构完全成立。

• 用RED指令替代普通原子操作无性能差异：CC7.5下全局内存的RED指令同样走绕过L1直达L2的路径，和普通全局原子的硬件执行流程完全一致，不会有性能差别。
• atomicAnd_block性能更差：这个接口的优化场景是block内所有线程集中对同一个内存地址做原子操作，和你当前“单warp内线程操作连续uint32地址、block内所有warp先处理完一组值再切下一行”的访问模式不匹配，接口自带的额外聚合指令开销反而会拉低性能。
• 无法测试redux指令、__shfl_sync表现差：redux是CC8.0才新增的warp级原子归约指令，CC7.5硬件没有对应的执行单元，确实无法运行；__shfl_sync仅支持同一warp内的线程交换数据，你的场景涉及跨warp的数据更新，没有共享内存作为中转的话shuffle根本无法完成跨warp交互，自然达不到预期效果。

• 放弃“L1自动缓存全局原子数据等效共享内存”的预期，这个硬件路径在Turing卡上不存在。
• 最优方案还是手动将当前待处理的行数据加载到共享内存，在共享内存上完成所有atomicAnd操作后再批量写回全局内存：共享内存原子的延迟比L2全局原子低一个数量级，__syncthreads()同步和数据拷贝的总开销，远低于直接发全局原子绕过L1访问L2的损耗，实际开发复杂度也远低于你预期。
• 如果确实不想手动管理共享内存，可以调整block内的warp数量，让同时发起的原子请求数和L2的原子处理通道数对齐，减少L2侧的请求排队，但这个方案的收益远低于共享内存方案。

内容的提问来源于stack exchange，提问作者David Wohlferd