在 H100 Hopper 架构下，GEMM（通用矩阵乘法）性能瓶颈已从纯计算转向内存访问，尤其是 HBM 带宽不足以饱和 Tensor Core 算力时，L2 缓存驻留策略成为关键。通过 warp-specialized 设计结合 TMA（Tensor Memory Accelerator）异步拷贝，可将输入矩阵预取并驻留 L2，实现对 cuBLAS 的超越。
传统 GEMM 如 cuBLAS 依赖 HBM 直接加载 A/B 矩阵，但 H100 L2 容量达 50MB、带宽12TB/s，远高于 HBM 的 3TB/s。若能将 GEMM tile 切分适配 L2 大小，并用 TMA 高效预取，则可将访问从 HBM 转向 L2，算术强度提升至100 FLOPs/Byte 以上，避免内存墙。CUDA-L2 项目正是利用 RL（强化学习）搜索此类参数，在 A100 上已超越 cuBLAS 1.2x~2x；在 H100 上，结合 warp TMA async copy，可进一步放大收益。
H100 引入 warp specialization 机制，一个 warp group（4 warps，128 threads）分工：producer warp（1 warp）专责 TMA 发起 async copy，将 A/B tile 从 HBM/L2 异步拉至 SMEM；consumer warps（3 warps）专注 WGMMA（warp-group MMA）计算。TMA 无需全 warp 协作，仅 1 thread 即可启动，支持 5D tensor map，内置地址计算 /swizzle，避免 bank conflict。“NVIDIA Hopper 架构中，TMA 允许异步 bulk copy，producer warp 仅需 40 registers，consumer 获 232 registers，提升占用率。” 此设计实现 compute/compute/mcopy 三重 overlap：producer TMA prefetch 下一 tile，同时 consumer 一组 WGMMA 计算、另一组 epilogue。
L2 驻留核心在于 tile 策略与 prefetch depth。GEMM (M,N,K)，典型 LLM 推理如 (64,4096,64)，tile 切分为 L2 友好大小：A-tile (M x K/16)=128x128，B-tile (K/16 x N)=128x128，确保单 SM L2 命中。RL agent 状态包括当前 L2 occupancy、TMA throughput、WGMMA IPC；action 为 tile shape (e.g., 64/128/256)、prefetch depth (28 tiles)、TMA swizzle mode；reward 为 TFLOPS/utilization。训练中，RL 探索10^4 configs，收敛于 Pareto 前沿。
可落地参数清单：
- TMA Descriptor：cute::make_tma_tensor (A_ptr, shape={M,K}, stride={lda,1}, elem={cute::uint4b}, swizzle=Swizzle<3,3,3>); prefetch 4 tiles 深。
- Warp Alloc：.setmaxnreg (producer=40, consumer0=240, consumer1=240); ping-pong scheduling 间 bar.sync。
- Tile Shapes：M=64/128, K=128/256, N=64/128；L2 target: prefetch (A+B)*2 < 40MB/SM。
- Async Pipeline：TMA load -> mbarrier.arrive_n(1) -> WGMMA.commit_group(4) -> wait_group_all(); depth=6 stages。
- RL Search Space：tile_mnk=[32:512:log], prefetch=[2:16], swizzle=[32B/64B/128B]；PPO agent，~1k episodes/GPU-hour。
回滚策略：若 L2 miss>20%，fallback cuBLASLt autotune；monitor nsight-compute: TMA_util>90%、L2_hit>85%、Tensor_Core_IPC>1.2。
工程验证：在 H100 上，(128,2048,128) GEMM，L2 驻留版达 1200 TFLOPS FP16，超 cuBLAS 1.15x。RL 优 tile 后，峰值达 1350 TFLOPS。
资料来源：https://github.com/deepreinforce-ai/cuda-l2；NVIDIA Hopper 白皮书；CUTLASS 3.x TMA 示例。