cuda-auto-tune

NCU 驱动的迭代式 kernel 优化（CUDA / CUTLASS / Triton / CuTe DSL）

门禁检查（在任何优化之前强制执行）

STOP — Do you have NCU profile data for this kernel?
  NO  → Go to Step 1. Do NOT touch any kernel code.
  YES → Go to Step 2.

硬性规则 — 违反任意一条都会使整个优化失效：

• 没有 NCU 数据时，绝不可更改 kernel 代码、启动配置或模板参数。
• 所有建议都必须引用具体的 NCU 指标值作为依据。
• 每次迭代至少必须覆盖：roofline、内存层级、warp 停顿、占用率。
• 优化手册必须与 kernel 的实现类型相匹配。
• 每次代码更改后，都要重新 profile 并使用 --diff 进行比较。
• 当改进趋于平缓或指标接近硬件上限时，停止迭代。

---

强制优化循环

┌─────────────────────────────────────────────────────────────────────┐
│  Step 1: Profile (NCU --set full)                                   │
│      ↓                                                              │
│  Step 2: Multi-dimensional analysis + identify kernel type          │
│      ↓                                                              │
│  Step 3: Apply type-specific playbook (one change per iteration)    │
│      ↓                                                              │
│  Step 4: Re-profile + diff → improved? → loop or stop               │
│      ↑                                           │                  │
│      └───────────────────────────────────────────┘                  │
└─────────────────────────────────────────────────────────────────────┘

---

步骤 1：使用 NCU 进行 profile（必需 — 没有数据就没有优化）

方案 A：Profiling 脚本（推荐）

# Native CUDA / CUTLASS binaries
bash cuda-auto-tune/scripts/ncu_profile.sh ./kernel report_v1

# Triton / Python
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_kernel.py" report_v1

# CuTe DSL / Python
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_cutedsl_kernel.py" report_v1

该脚本会收集 --set full → 导出 CSV → 运行深度分析 → 生成报告。

方案 B：手动 profiling

ncu --set full -o report_v1 --target-processes all ./your_kernel
ncu --import report_v1.ncu-rep --page raw --csv > report_v1.csv
python3 cuda-auto-tune/scripts/ncu_analyse.py report_v1.csv

Kernel 名称过滤器（减少噪声）

# CUTLASS only
ncu --set full -o report_v1 --target-processes all \
    --kernel-name "cutlass_\|sm90_\|ampere_" ./cutlass_program

# Triton only
ncu --set full -o report_v1 --target-processes all \
    --kernel-name "triton_" "python triton_kernel.py"

# CuTe DSL (kernel name often generic — use --type override in analysis)
python3 cuda-auto-tune/scripts/ncu_analyse.py report_v1.csv --type cutedsl

预期输出

ncu_reports/
├── report_v1.ncu-rep           # Full binary report
├── report_v1.csv               # Raw metrics CSV
├── report_v1_analysis.md       # Deep analysis report
└── report_v1_summary.txt       # Per-kernel summary

---

步骤 2：多维度分析

2.1 识别实现类型

根据 NCU 的 "Function Name" 和源代码上下文确定 kernel 类型：

类型	检测信号
原生 CUDA	没有库前缀；手写的 __global__ 函数
CUTLASS	cutlass_ 前缀、smXX_xmma_、包含 tensorop 或 cutlass
Triton	triton_ 前缀、包含 triton、带编码的后缀（例如 _0d1d...e）
CuTe DSL	来自 @cute.kernel 的通用名称；通过源代码导入（cutlass.cute、cute.compile）或 --type cutedsl 确认
库	cublas*、cudnn* — 仅作基线/参考，不可优化

2.2 通用诊断（所有 kernel 类型 — 始终运行）

维度	关键 NCU 指标	输出
Roofline	SM 吞吐量、内存吞吐量	compute-bound / memory-bound / latency-bound / balanced
内存层级	L1/L2 命中率、合并访问比率、DRAM 吞吐量	缓存效率 + 带宽子瓶颈（DRAM/L2/L1）
Warp 停顿	PC 采样停顿原因（long_scoreboard、wait、barrier……）	带百分比的 top 停顿原因
指令构成	pipe FMA/ALU/LSU/Tensor 利用率	流水线不均衡、Tensor Core 使用情况
占用率	active warps %、限制因素分解（register/smem/warp/block）	限制因素 + 寄存器数量 + smem 大小
内存隐患	bank conflicts、寄存器溢出（local store sectors）	严重程度与根本原因
发散	平均执行线程数 vs 平均活跃线程数（true）	发散百分比

2.3 类型特定的关注点

类型	关键关注领域
原生 CUDA	启动配置（block 大小、grid）、内存访问模式、异步拷贝（cp.async/TMA）、Tensor Core 机会
CUTLASS	ThreadblockShape、WarpShape、stages、对齐、调度策略、epilogue 融合、CTA swizzle
Triton	num_warps、num_stages、BLOCK_* 大小、编译器提示（tl.multiple_of、tl.max_contiguous）、tl.dot 配置
CuTe DSL	threads_per_cta、elems_per_thread、CopyAtom（num_bits_per_copy）、tiled_copy 布局、smem 暂存、cta_reduce 模式

2.4 瓶颈分类决策树

SM% > MEM% + 20  →  COMPUTE_BOUND
MEM% > SM% + 20  →  MEMORY_BOUND
  ├─ DRAM throughput > 70%        → DRAM-Bound (near HBM ceiling)
  ├─ L2 hit < 50%, DRAM > 40%    → DRAM-Bound (L2 miss driven)
  ├─ L1 hit < 20%, L2 hit >= 50% → L2-Bound
  └─ L1 hit < 20%                → L1-Bound
SM% < 40 AND MEM% < 40           →  LATENCY_BOUND
SM% > 60 AND MEM% > 60           →  BALANCED (near peak)

2.5 结论模板（每次分析后必需）

=== Conclusion ===
Kernel:    {kernel_name}
Type:      {Native CUDA | CUTLASS | Triton | CuTe DSL}
Arch:      SM_{arch}
Overall:   {COMPUTE_BOUND | MEMORY_BOUND | LATENCY_BOUND | BALANCED}
Duration:  {duration_us} us
Roofline:  SM {sm}%, MEM {mem}%, DRAM {dram}%
Occupancy: {occ}% (theoretical: {theo}%), limited by {limiter}
Regs/Thread: {regs}, Smem/Block: {smem} KB

Findings (sorted by severity):
  [CRITICAL] {finding}: {NCU evidence with numbers} -> {specific action}
  [WARNING]  {finding}: {NCU evidence with numbers} -> {specific action}
  [INFO]     {finding}: {NCU evidence with numbers}

Optimization priorities:
  1. {highest_priority} (expected gain: Nx, evidence: {metric}={value})
  2. {second_priority}  (expected gain: Nx, evidence: {metric}={value})
  3. {third_priority}   (expected gain: Nx, evidence: {metric}={value})

---

步骤 3：应用类型特定的手册

不要仅凭直觉修改。每项更改都必须直接针对一条 NCU 发现。 每次迭代只应用一项更改，然后重新 profile（步骤 4）。

---

3.1 手册：原生 CUDA

3.1.1 启动配置

NCU 发现	操作	代码模式
占用率 < 50%，线程块大小 < 128	将线程块大小增至 128–256	kernel<<<grid, 256>>>
寄存器是占用率限制因素	通过 __launch_bounds__ 限制寄存器	__global__ void __launch_bounds__(256, 2) kernel()
网格过小（< SM 数量）	确保有足够的线程块覆盖全部 SM	grid = (N + block - 1) / block，且 N 足够大
占用率低，线程块是限制因素	减小线程块大小，使每个 SM 容纳更多线程块	尝试用 128 替代 256

3.1.2 内存访问优化

NCU 发现	操作	代码模式
加载合并比率 > 8	确保 warp 连续寻址，AoS→SoA	data[threadIdx.x + blockIdx.x * blockDim.x]
存储合并比率 > 8	对分散写入使用共享内存中转	先写入 smem，再合并写回
L1 命中率 < 20%	对频繁复用的数据使用 __shared__	通过 __syncthreads() 分块载入共享内存
L2 命中率 < 50%	使用 L2 持久化提示（Ampere+）	对热点数据范围使用 cudaAccessPolicyWindow
DRAM 吞吐 > 80%	减少数据搬运：混合精度、压缩	对带宽敏感的操作使用 half / __nv_bfloat16
Bank 冲突 > 100K	对共享内存填充或重排布局	__shared__ float smem[32][33];（填充 +1）
寄存器溢出 > 0	减少每线程状态，使用 __launch_bounds__	简化累加器，拆分为子 kernel

3.1.3 延迟隐藏与流水线

NCU 发现	操作	代码模式
stall_long_scoreboard > 30%（SM>=80）	使用 cp.async + 双缓冲	__pipeline_memcpy_async(&smem, &gmem, size)
stall_long_scoreboard > 30%（SM>=90）	使用 TMA 进行批量异步传输	cute::copy(tma_load, ...) 或 CuTe TMA atoms
stall_barrier > 25%	降低同步频率，使用 warp 原语	__shfl_sync()、cooperative_groups
stall_wait > 30%，long_scoreboard < 15%	流水线缓冲过多，减少深度	移除一级缓冲
stall_math_pipe_throttle > 20%	计算已饱和（正面信号）	考虑使用 Tensor Core 或减少 FLOPs

3.1.4 Tensor Core 利用率

NCU 发现	操作
pipe_tensor < 5%，FP16/BF16 工作负载且呈 GEMM 类模式	使用 WMMA（wmma::mma_sync）或内联 PTX（mma.sync）
pipe_tensor < 5%，但数据为 FP32	通过 wmma::mma_sync 配合 nvcuda::wmma::precision::tf32 使用 TF32 路径
pipe_fma_fp16 > 10%，pipe_tensor < 5%	从标量 FP16 FMA 切换到 Tensor Core 路径

3.1.5 向量化内存访问

// NCU evidence: coalescing ratio > 4 for 32-bit loads
// Before: scalar loads
float val = input[idx];

// After: vectorized 128-bit load (4x float)
float4 val = reinterpret_cast<const float4*>(input)[idx / 4];

---

3.2 实战手册：CUTLASS

3.2.1 Kernel 配置解析

CUTLASS kernel 名称编码了配置信息。提取以下内容：

• 架构：sm80_、sm90_、ampere_、hopper_
• 计算类型：tensorop 与 simt
• Tile 形状：128x128x32、256x128x64
• 流水线级数：末尾的 x3、x5
• 对齐：align8
• 调度（3.x）：WarpSpecialized、WarpSpecializedCooperative、WarpSpecializedPingpong

3.2.2 Tile 形状与占用率

NCU 发现	操作
占用率 < 40%，smem 是限制因素	减小 ThreadblockShape（如 256x128→128x128）或减少级数
占用率 < 40%，寄存器是限制因素	使用更小的 WarpShape（如 64x64→32x32）以减少每线程寄存器数
SM 吞吐 < 30%，网格较小	增大 ThreadblockShape，使每个 CTA 处理更多元素
SM 吞吐 > 80%，MEM < 40%	已是计算受限；增加流水线级数以获得更多重叠

3.2.3 流水线级数

NCU 发现	操作
stall_long_scoreboard > 30%	增加级数（Ampere：3→5，Hopper：2→3）
stall_wait > 30%，long_scoreboard < 15%	流水线缓冲过多；减少级数以节省 smem
Smem 受限 + 级数 > 3	减少级数以释放 smem，提高占用率

3.2.4 对齐与向量化

NCU 发现	操作
Load coalescing > 4，alignment < 8	将 CUTLASS alignment 提升至 8（128 字节）；将矩阵 leading dims 填充至 alignment 的倍数
使用了 SIMT 路径但数据支持 TensorOp	切换到 tensorop CUTLASS 配置（2–8 倍加速）
已配置 TensorOp 但 pipe_tensor < 5%	检查对齐要求 —— LD 必须是 InstructionShape::kK 的倍数

3.2.5 调度与架构

NCU 发现	操作
SM>=90 上使用 CUTLASS 2.x	升级到 CUTLASS 3.x，使用 WarpSpecialized + TMA（1.2–1.5 倍收益）
大型 GEMM 上 L2 命中率 < 50%	添加 ThreadblockSwizzle（2.x：GemmIdentityThreadblockSwizzle<N>，3.x：StreamK 或 tile swizzle）
Hopper 上 stall_long_scoreboard > 30%	切换到使用 TMA 加载的 WarpSpecializedCooperative schedule

3.2.6 Epilogue 融合

NCU 发现	操作
多个 CUTLASS kernel 连续执行（例如 GEMM + bias + activation）	通过 CUTLASS epilogue visitor tree 融合为单个 kernel
DRAM 流量高（读+写 GB 超出预期）	将 GEMM 后的操作移入 epilogue，以消除中间张量

---

3.3 Playbook：Triton

3.3.1 Kernel 分类

Triton kernel 子类型（来自 kernel 名称）：

• triton_poi_：Inductor pointwise（自动生成）
• triton_red_：Inductor reduction（自动生成）
• triton_per_：Inductor persistent reduction（自动生成）
• 自定义 @triton.jit：手写 kernel（完全可调优）

Inductor 生成的 kernel：在 PyTorch 层面优化（torch._inductor.config），或者如果是热点路径，则重写为自定义 @triton.jit。

3.3.2 num_warps 调优

NCU 发现	操作
Registers >= 128，num_warps >= 8	CRITICAL：减少 num_warps（尝试 4 或 2）
Registers >= 64，num_warps >= 8	将 num_warps 减少到 4
Occupancy < 40%，受寄存器限制	减少 num_warps 和/或减少 BLOCK_* tile 尺寸
SM 吞吐 < 30%，warp 数量少	增加 num_warps 以改善延迟隐藏

3.3.3 num_stages 调优

NCU 发现	操作
stall_long_scoreboard > 30%	增加 num_stages（Ampere 上 2→3→4，Hopper 上 2→3）
stall_wait > 30%，long_scoreboard < 15%	减少 num_stages（缓冲区过量）或增加 tile 工作量
Smem 是 occupancy 限制因素	减少 num_stages（每个 stage 会使 smem 缓冲区翻倍）
在 Hopper 上且 long_scoreboard 高	同时考虑使用 tl.make_block_ptr() 进行基于 TMA 的加载

3.3.4 BLOCK_* tile 尺寸调优

NCU 发现	操作
寄存器压力高	减少 BLOCK_M、BLOCK_N 或 BLOCK_K
SM 吞吐低，存在计算受限机会	增加 BLOCK_M/BLOCK_N，以增加每个 tile 的计算量
DRAM 带宽接近上限	增加 BLOCK_K，以在写回前复用更多数据

3.3.5 内存访问优化

NCU 发现	操作	代码模式
Load coalescing > 8	添加 stride 提示	tl.multiple_of(stride, 16) 和 tl.max_contiguous(offsets, BLOCK)
转置输入上未合并访问	使用结构化指针	tl.make_block_ptr(base, shape, strides, offsets, block_shape, order)
L1 命中率低	验证访问模式的连续性	确保最内层维度 stride == 1

3.3.6 Tensor Core 利用率

NCU 发现	操作
pipe_tensor < 5%，kernel 使用 tl.dot	1）对 fp32 使用 allow_tf32=True；2）BLOCK_K 为 16 的倍数；3）检查 dtype 为 fp16/bf16/tf32/fp8
pipe_tensor < 5%，代码中没有 tl.dot	缺少类 GEMM 模式 —— 重构以使用 tl.dot

3.3.7 Triton autotune 集成

@triton.autotune(
    configs=[
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 128, 'BLOCK_K': 32}, num_warps=4, num_stages=3),
        triton.Config({'BLOCK_M': 64,  'BLOCK_N': 64,  'BLOCK_K': 64}, num_warps=4, num_stages=4),
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 64,  'BLOCK_K': 32}, num_warps=8, num_stages=3),
    ],
    key=['M', 'N', 'K'],
)
@triton.jit
def kernel(...):
    ...

当 NCU 揭示瓶颈时，应将 autotune 配置缩小到有希望的区域，而不是盲目搜索。

---

3.4 Playbook：CuTe DSL

3.4.1 关键调优参数

参数	作用	典型范围
threads_per_cta	每个 CTA 的 warp 数；影响占用率、屏障开销、reduce 开销	128–512
elems_per_thread	每个线程处理的元素数；影响寄存器压力、数据复用	4–32
num_bits_per_copy	CopyAtom 宽度；影响向量化加载/存储带宽	32, 64, 128
Smem 暂存缓冲区	流水线深度 × tile 大小；影响 smem 占用	为提升占用率应尽量最小化

3.4.2 占用率优化

NCU 发现	操作
占用率 < 40%，寄存器是限制因素	减少 elems_per_thread 或减少 threads_per_cta；在 cute.compile() 中添加 --maxrregcount=128
占用率 < 40%，smem 是限制因素	减少 threads_per_cta（更少的 warp → 更小的 reduce 缓冲区）或减少暂存缓冲区数量
寄存器 >= 128，warp >= 8	CRITICAL：将 threads_per_cta 减少到 128 或 256

3.4.3 内存访问（TiledCopy）

NCU 发现	操作
加载合并度 > 8	1) 将 num_bits_per_copy 提高到 128；2) 验证 t_layout 沿连续地址分布线程；3) 确保 from_dlpack() 使用 assumed_align=16
stall_long_scoreboard > 30%	1) 将 num_bits_per_copy 提高到 128；2) 增加 elems_per_thread 以实现更多复用；3) 在 SM>=80 上使用 CpAsyncOp copy atom；4) 添加双缓冲
stall_wait > 30%，long_scoreboard < 15%	流水线缓冲区过多；增加 elems_per_thread 以在每个阶段进行更多计算，或减少流水线深度

3.4.4 同步与归约

NCU 发现	操作
stall_barrier > 25%	1) 减少 threads_per_cta（屏障处更少的 warp）；2) 将第二个 sync_threads 替换为 shuffle 广播（如果 warp <= 32）；3) 合并多个 cta_reduce 调用
高屏障 + 小归约	对于少量元素，使用仅 warp 归约而不使用 smem
每次迭代多次 sync_threads	最小化同步点；使用异步流水线 commit/wait 模式

3.4.5 线程分歧

NCU 发现	操作
分歧 > 20%	调整 threads_per_cta * elems_per_thread 以紧密匹配问题维度 N，减少被谓词关闭的线程
谓词化拷贝显示高分歧	确保 N 能被 threads_per_cta * elems_per_thread 整除，或使用尾部处理策略

3.4.6 计算优化

NCU 发现	操作
pipe_tensor < 5%，类 FP16 GEMM 操作	使用 cute.make_mma_atom() 搭配 MmaOp 走 Tensor Core 路径
pipe_fma 高但 pipe_tensor 低（如 RMSNorm/LayerNorm 等非 GEMM 操作）	Tensor Core 不适用于归约 — 应转而专注于内存和屏障优化

3.4.7 重新分析前的缓存失效

CuTe DSL 通过 JIT 将 Python 编译为 CUDA。代码更改后：

# Clear compilation cache to ensure re-compilation
rm -rf __pycache__/ .cache/ /tmp/cutlass_cute_cache/
# Then re-profile
bash cuda-auto-tune/scripts/ncu_profile.sh "python your_cutedsl_kernel.py" report_v2

---

步骤 4：重新分析和验证（每次更改后必需）

4.1 重新分析

# Clear JIT caches first
rm -rf ~/.triton/cache            # Triton
rm -rf __pycache__/ .cache/       # CuTe DSL

# Profile updated version
bash cuda-auto-tune/scripts/ncu_profile.sh ./kernel_v2 report_v2
# or
bash cuda-auto-tune/scripts/ncu_profile.sh "python kernel_v2.py" report_v2

4.2 与基线对比

python3 cuda-auto-tune/scripts/ncu_analyse.py ncu_reports/report_v2.csv --diff ncu_reports/report_v1.csv

4.3 验证清单

检查项	标准
耗时有改善？	gpu__time_duration.sum 降低
目标瓶颈有改善？	触发更改的特定指标有改善
无新瓶颈？	差异报告中无新的 CRITICAL 发现
达到硬件上限？	SM 吞吐量 > 80% 或 DRAM 吞吐量 > 85% 表示接近峰值

4.4 迭代日志模板

跟踪每次迭代以便问责：

=== Iteration {N} ===
Change:  {what was changed and why}
NCU evidence: {metric}={before_value} -> {finding}
Report: report_v{N}.csv

Result:
  Duration: {before} us -> {after} us ({delta}%)
  Target metric: {metric}={before} -> {after}
  New findings: {any new issues introduced}

Decision: {CONTINUE to next bottleneck | STOP — at ceiling | ROLLBACK — regression}

---

快速参考：高信号 NCU 指标

指标	NCU 键
时长	gpu__time_duration.sum [us]
SM 吞吐量	sm__throughput.avg.pct_of_peak_sustained_elapsed [%]
内存吞吐量	gpu__compute_memory_throughput.avg.pct_of_peak_sustained_elapsed [%]
DRAM 吞吐量	gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed [%]
L1 命中率	l1tex__t_sector_hit_rate.pct [%]
L2 命中率	lts__t_sector_hit_rate.pct [%]
加载合并	l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum / l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum
Bank 冲突	l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum
寄存器溢出	l1tex__t_sectors_pipe_lsu_mem_local_op_st.sum [sector]
占用率	sm__warps_active.avg.pct_of_peak_sustained_active [%]
Warp 可调度性	smsp__warps_eligible.avg.per_cycle_active [warp]
寄存器/线程	launch__registers_per_thread [register/thread]
共享内存/块	launch__shared_mem_per_block [Kbyte/block]

---

摘要

本技能强制执行严格的 profile → analyze → change → verify 循环。 没有 NCU 数据 = 不进行优化。没有指标证据 = 不修改代码。 每种 kernel 类型（Native CUDA / CUTLASS / Triton / CuTe DSL）都有专门的 playbook， 其中包含 NCU 指标到操作的映射。每次更改都会被跟踪，并通过重新 profile 进行验证。