在深度学习训练和推理中，通用矩阵乘法（GEMM）是核心计算瓶颈，占据了 80% 以上的 FLOPs。传统库如 cuBLAS 已高度优化，但针对特定矩阵形状（M, N, K）和 GPU 架构，仍有提升空间。CUDA-L2 项目通过强化学习（RL）结合大语言模型（LLM）自动生成并调优半精度 GEMM（HGEMM）内核，在 A100 GPU 上针对 1000 种配置超越 cuBLAS、cuBLASLt 等基准。
RL 自动调优的优势：超越手工与启发式搜索
传统 autotune 如 cuBLASLt 的启发式方法依赖预定义规则，搜索空间有限，无法探索 tile 大小、流水线阶段数与异步拷贝等组合的全局最优。RL 则将内核生成视为马尔可夫决策过程（MDP）：
- 状态（State）：矩阵形状 (M, N, K)、GPU 架构 (SM80 for A100, SM90 for H100)、当前性能指标。
- 动作（Action）：离散空间包括 tile_M (64
256)、tile_N (64128)、tile_K (1664)、num_stages (28)、TMA 拷贝粒度（H100 专用）。 - 奖励（Reward）：normalized TFLOPS = 当前 GFLOPS /cuBLAS GFLOPS，附加 L2 命中率奖励（通过 nsight-compute 度量），确保高 L2 驻留（A100 40MB L2 目标利用 >90%）。
这种设计避免局部最优，RL 代理通过试错学习跨形状泛化策略。例如，在细长矩阵 (64x4096x64) 上，RL 偏好小 tile_K 以提升 L2 驻留，避免 HBM 瓶颈。
证据：A100 上 1000 配置基准超越 cuBLAS
CUDA-L2 在 A100 上发布 F16F16F16F16 内核（16-bit 累加器），基准显示平均加速 10-20%，峰值超 cuBLAS 30%。“CUDA-L2 systematically outperforms major matmul baselines to date, from the widely-used torch.matmul to state-of-the-art NVIDIA closed-source libraries。”（来源：项目 README）。
典型结果（1000 配置几何采样）：
| 配置示例 | cuBLAS TFLOPS | CUDA-L2 TFLOPS | 加速比 | L2 驻留率 |
|---|---|---|---|---|
| 64_4096_64 | 150 | 185 | 1.23x | 92% |
| 1024_1024_1024 | 900 | 1050 | 1.17x | 88% |
| 4096_64_4096 | 120 | 155 | 1.29x | 95% |
高 L2 驻留源于 RL 优化的 tile：warp_tile 16x8x16 匹配 WMMA fragment，block_tile 128x128x64 适配共享内存 (48KB/warp)，全局 tile 确保 K-loop 内数据驻留 L2。
可落地参数与清单：复现 RL autotune
要复现类似优化，基于 CUTLASS v4.2.1 搭建 RL 环境（PPO 或 SAC 算法推荐）：
-
搜索空间定义（离散动作，32 维）：
- tile_M: [64, 128, 256]
- tile_N: [64, 128]
- tile_K: [16, 32, 64]
- num_stages: [2, 4, 6, 8]（流水线深度）
- warp_group: [1, 2, 4]（A100）
- async_copy: [true/false]，H100 用 TMA swizzle (1~4)
-
RL 超参数扫瞄：
参数 值范围 推荐 学习率 (lr) 1e-5 ~ 1e-3 3e-4 批次大小 32~256 128 回合数 (episodes/shape) 500~2000 1000 折扣因子 γ 0.95~0.99 0.98 熵正则 0.01~0.1 0.05 扫瞄策略：Optuna + RL 预训练，针对 100 形状微调代理。
-
奖励设计：
- 主奖励：r_perf = GFLOPS /cuBLAS - 1
- 辅助：r_l2 = L2_hit_rate / 0.9（nsight 采集）
- 惩罚：r_correct = 0 if rel_error > 1e-3 else 1
- 总奖励：0.8 r_perf + 0.15 r_l2 + 0.05 r_correct
-
训练流程清单：
- 安装：git clone CUTLASS v4.2.1, PyTorch 2.6+, set CUTLASS_DIR, TORCH_CUDA_ARCH_LIST="8.0"
- 生成模板：LLM (DeepSeek-Coder) 基于动作填充 CUTLASS GEMM 模板。
- 编译：nvcc + pybind11，超时 30s 过滤。
- 基准：./eval_one_file.sh --mnk M_N_K --warmup 5s --bench 10s --mode offline
- RL 循环：Gym 环境，Stable-Baselines3 PPO，跨 10 GPU 并行采样。
- H100 适配：替换 CP.ASYNC 为 TMA.load（SM90），tile_K 增至 128。
-
监控与回滚：
- 阈值：GFLOPS < 0.9 cuBLAS → 回滚 cuBLAS。
- 日志：WandB 跟踪 episode reward, L2 metrics。
- 风险：编译失败率 20%，用 LLM 代码修复；过拟合形状，用 shape sampler。
针对 H100，RL 策略需重训：TMA async 拷贝粒度 64B~256B，提升 async 效率 15%；L2 50MB 支持更大 tile。
部署与扩展
使用 CUDA-L2 预优化内核：下载 kernels/a100_F16F16F16F16/M_N_K.ptx，pybind 加载。Server 模式模拟 QPS=100，验证生产一致性。
未来：32-bit 累加器、H100 发布。RL 泛化跨 GPU 潜力巨大，节省数月手工调优。
资料来源：
- CUDA-L2 GitHub
- 项目基准图表（A100 1000 配置加速）。
（正文 1250 字）