| name | ncu-report |
| description | Profile CUDA kernels with Nsight Compute on SM89 (RTX 4090/4070). Use when the user asks to profile a kernel, analyze its performance, diagnose bottlenecks, read an ncu report, or write an optimization plan — including variants like "profile 一下", "为什么慢", "ncu 报告", "性能分析". |
Skill: CUDA Kernel Profiling (SM89 / Nsight Compute)
When to use: user asks to profile a CUDA kernel, analyze its performance, find bottlenecks, or write an optimization plan based on Nsight Compute data. Triggers include: "profile X", "为什么这个 kernel 慢", "ncu report", "下一步怎么优化", "帮我看一下这份 ncu 报告", "性能分析".
Target hardware: NVIDIA RTX 4090/4070 (SM89, CC 8.9, 128 SM, ~1 TB/s DRAM BW, 4th gen Tensor Core with FP8/FP16/BF16/TF32/INT8).
Golden Rule
Profile → Diagnose → Plan, in that order. Never guess.
Most under-performing CUDA kernels are under-performing for exactly one reason that ncu can tell you in 10 seconds. Don't invent hypotheses before you have the report. Don't start coding a fix before you've matched the observed pattern to a known diagnosis.
Quick Start — "Profile this kernel"
Step 1: Build (确保有 line-info)
本项目的 CMake target 已默认带 --generate-line-info。直接编译：
cmake -S . -B build -DTINY_CUTLASS_BUILD_<FAMILY>=ON
cmake --build build --config Release --target <target_name>
Step 2: Verify Correctness (必须先过 correctness gate)
build\csrc\<family>\Release\<target>.exe --reference-check=true [small shape args]
Step 3: Profile (一条命令搞定)
python scripts/profile/run_ncu.py ^
--run-name <descriptive_name> ^
--exe build\csrc\<family>\Release\<target>.exe ^
--kernel-regex "<kernel_name_regex>" ^
--args "<exe_args>" ^
--set both --count 1
这会自动：
- 创建
profile/<run_name>/ 目录结构
- 执行
ncu --set full + PmSampling
- 执行
ncu --set source + SourceCounters
- 导出 details page + CSV
- 调用
analyze_report.py / extract_stalls.py / plot_timeline.py
Step 4: Diagnose
按六维分析法检查 profile/<run_name>/analysis/ 下的产出：
metrics_key_default.txt — headline numbers + occupancy + stallsdetails_default.txt — NCU 规则引擎建议（先看这个，每条带 Est. Speedup: X%）stall_hotspots_default.txt — per-line stall 热点pm_timeline_plots.txt — 时序图（揭示 tail effect / pipeline bubble）
Step 5: Write REPORT.md
在 profile/<run_name>/REPORT.md 中写诊断报告。用 reference/report_template.md 的模板。
File Index
工具脚本 (scripts/profile/)
| File | Purpose |
|---|
run_ncu.py | 统一入口 — 创建目录 → ncu 收集 → 导出 → 分析，一条命令搞定 |
ncu_utils.py | 核心工具：加载 .ncu-rep、安全 metric 读取、SM89 key metrics 列表 |
analyze_report.py | 提取 key metrics + 多报告对比 |
extract_stalls.py | Per-line stall hotspots（需要 source-level 报告） |
plot_timeline.py | ASCII PM sampling 时序图 |
Helpers (scripts/profile/helpers/)
| File | Purpose |
|---|
harness_template.cu | 独立 profiling harness 模板（当不想走 CMake 时使用） |
safetensors_loader.h | Header-only safetensors 读取器（加载真实 workload tensor） |
README.md | 使用说明 |
Reference Docs (scripts/profile/reference/)
| File | Purpose |
|---|
directory_layout.md | 先读 — 目录命名规则 |
workflow.md | 端到端 profiling 流程（Phase 0-6） |
harness_guide.md | 何时/如何构建独立 harness |
collection.md | ncu 命令 recipes（SM89 + Windows） |
python_api.md | ncu_report Python API 详细参考 |
analysis_dimensions.md | 六维分析法 |
diagnosis_playbook.md | Signal → Diagnosis → Fix（14 patterns A-N） |
report_template.md | REPORT.md 模板与风格指南 |
sm89_metrics.md | SM89 metric name 完整参考 |
common_issues.md | 权限、PM Sampling、匹配失败等排查 |
Critical Lessons
-
先看 details_<tag>.txt。 NCU 规则引擎的 Est. Speedup: X% 通常直接指向答案。
-
编译必须带 --generate-line-info。 本项目 CMake 已默认开启。如果 source view 空白 = 没有 line info。
-
PM sampling 是唯一能看到 tail effect 的方式。 静态 metric 做全 kernel 平均；只有时序图能揭示利用率随时间的形状。
-
不要委托理解。 自己跑 profile，打开报告，引用具体 metric 值。不要写"profile 显示是 memory-bound" — 而是说明哪两三个 metric 值支撑了这个结论。
-
Sectors/request = 4 是理想值。 > 8 = 严重 uncoalesced。用 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum / l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 计算。
-
local_ld > 0 = register spill。 这是 DRAM 速度的溢出。用 __launch_bounds__ 或拆 kernel。
Driver Script (Smoke Test)
验证整个 profiling 工具链工作正常：
rem 如果有已编译的 target 可用：
python scripts/profile/run_ncu.py ^
--run-name smoke_test ^
--exe build\csrc\flash-attention\Release\naive_attention.exe ^
--kernel-regex ".*" ^
--args "--batch_size=1 --head_number=1 --head_size=32 --head_size_v=32 --seq_length=32 --seq_length_kv=32 --iterations=1 --reference-check=false" ^
--set full --count 1
rem 检查产出
dir profile\smoke_test\analysis\
type profile\smoke_test\analysis\metrics_key_default.txt
如果 ncu_report import 失败：
set PYTHONPATH=%PYTHONPATH%;C:\Program Files\NVIDIA Corporation\Nsight Compute 2025.2.0\extras\python
Integration with cutlass-kernel Skill
本 skill 与 cutlass-kernel skill 配合使用：
cutlass-kernel 定义 kernel 开发的 build → verify → bench 工作流ncu-report 在 bench 之后提供 深度性能分析 能力
典型流程：
cutlass-kernel workflow: build → verify → bench → 发现 kernel 比 cuBLAS 慢ncu-report workflow: profile → diagnose → 写 REPORT.md → 得到优化方向- 回到
cutlass-kernel: 按优化方向实现新 variant → verify → bench → 确认提升
