Menu
Automated NCU (Nsight Compute) profiling workflow with full metrics collection and persistent storage
Install with Codex or Claude Copy this prompt, paste it into Codex, Claude, or another assistant, and let it review the skill page and install it for you.
Based on SOC occupation classification
| name | ncu-cuda-profiling |
| description | Automated NCU (Nsight Compute) profiling workflow with full metrics collection and persistent storage |
| version | 1.0.0 |
| author | maxiaosong1124 |
| tags | ["cuda","profiling","ncu","performance","optimization"] |
本 Skill 提供完整的自动化 NCU 性能分析流程，支持全量指标采集和持久化存储。
# 使用 --set full 采集所有指标，并持久化保存
ncu --set full \
-o <report_name> \
--target-processes all \
./your_kernel
# 示例
ncu --set full -o matmul_analysis --target-processes all ./matmul0_perf
# 自动生成:
# - matmul_analysis.ncu-rep (NCU 报告文件)
# - matmul_analysis.csv (CSV 格式指标)
# 从已保存的报告提取关键指标 (无需重新运行 kernel)
ncu --import matmul_analysis.ncu-rep --print-summary per-kernel
# 导出为 CSV
ncu --import matmul_analysis.ncu-rep --page raw --csv > metrics.csv
当用户提供 NCU 数据时，AI 按以下流程处理：
情况 A: 用户提供了 .ncu-rep 文件
# 直接导入已有报告
ncu --import <file.ncu-rep> --print-summary per-kernel
情况 B: 用户需要新分析
# 完整采集并持久化
ncu --set full -o <report_name> --target-processes all ./kernel
情况 C: 用户提供了截图/文本
AI 会自动保存分析数据到项目目录：
project_root/
├── ncu_reports/ # NCU 报告目录
│ ├── matmul_analysis.ncu-rep # 完整报告
│ ├── matmul_analysis.csv # CSV 指标
│ └── matmul_analysis.md # AI 分析报告
└── ...
使用决策引擎自动分析：
def auto_diagnose(metrics):
roofline = metrics.get('roofline_ratio', 0)
dram = metrics.get('dram_throughput', 0)
l1tex = metrics.get('l1tex_throughput', 0)
sm_busy = metrics.get('sm_busy', 0)
occupancy = metrics.get('occupancy', 0)
if roofline < 30:
if dram > 70:
return "DRAM_MEMORY_BOUND"
elif l1tex > 80 and dram < 30:
return "L1_PRESSURE_BOUND"
else:
return "LATENCY_BOUND"
elif roofline > 60:
if sm_busy > 80:
return "COMPUTE_BOUND"
else:
return "OCCUPANCY_BOUND"
else:
return "MIXED_BOUND"
# NCU 性能分析报告
## 📁 报告信息
- **Kernel**: {kernel_name}
- **采集时间**: {timestamp}
- **报告文件**: {report_file}
- **原始数据**: {csv_file}
## 📈 执行摘要
| 项目 | 数值 |
|------|------|
| **主要瓶颈** | {bottleneck_type} |
| **置信度** | {confidence} |
| **性能** | {performance} GFLOPS |
| **优化潜力** | {potential}x |
## 📊 关键指标
### 性能指标
| 指标 | 数值 | 健康阈值 | 状态 |
|------|------|----------|------|
| Roofline 性能比 | {roofline}% | > 60% | {status} |
| SM Busy | {sm_busy}% | > 70% | {status} |
| Occupancy | {occupancy}% | > 50% | {status} |
### 内存指标
| 指标 | 数值 | 健康阈值 | 状态 |
|------|------|----------|------|
| DRAM Throughput | {dram}% | < 50% | {status} |
| L1/TEX Throughput | {l1tex}% | < 80% | {status} |
| L2 Throughput | {l2}% | < 80% | {status} |
## 🔍 诊断详情
**瓶颈类型**: {bottleneck_type}
**判断依据**:
- {reason_1}
- {reason_2}
## 💡 优化建议
### 高优先级
{high_priority_suggestions}
## 🛠️ 下一步操作
### 建议的 NCU 命令
```bash
# 优化后重新采集
ncu --set full -o {report_name}_optimized --target-processes all ./kernel_optimized
---
## 🔧 工具使用说明
### 完整采集 (推荐)
```bash
# 采集所有指标并保存
ncu --set full -o my_analysis --target-processes all ./kernel
# 参数说明:
# --set full # 采集完整指标集
# -o my_analysis # 输出文件名 (生成 my_analysis.ncu-rep)
# --target-processes all # 监控所有进程
# 从已有报告提取特定指标
ncu --import my_analysis.ncu-rep --print-summary per-kernel
# 导出为 CSV 便于分析
ncu --import my_analysis.ncu-rep --page raw --csv > metrics.csv
使用提供的自动化脚本：
cd examples/
# 全自动分析
./auto_profile.sh ./kernel report_name
# Python 分析器
python ncu_analyzer.py --import report_name.ncu-rep
IF dram_throughput > 70% AND roofline < 30%:
诊断: DRAM_MEMORY_BOUND (置信度: HIGH)
优化策略:
1. Block Tiling (共享内存缓存)
2. Vectorized Load (float4)
3. Prefetching (数据预取)
IF l1tex_throughput > 80% AND dram_throughput < 30%:
诊断: L1_PRESSURE_BOUND (置信度: HIGH)
优化策略:
1. Shared Memory Padding
2. Data Transpose
3. Fragment Caching
IF sm_busy < 50% AND occupancy > 60%:
诊断: LATENCY_BOUND (置信度: HIGH)
优化策略:
1. Double Buffering
2. Instruction-level Parallelism
3. Loop Unrolling
IF roofline > 60% AND sm_busy > 80%:
诊断: COMPUTE_BOUND (置信度: HIGH)
优化策略:
1. Use FMA instructions
2. Reduce precision (FP32 -> FP16/TF32)
3. Tensor Cores
IF occupancy < 30% AND sm_busy > 70%:
诊断: OCCUPANCY_BOUND (置信度: HIGH)
优化策略:
1. Reduce register usage
2. Adjust block size
3. Use __launch_bounds__
| 瓶颈类型 | 立即行动 | 代码示例 | 预期收益 |
|---|---|---|---|
| DRAM_MEMORY_BOUND | Block Tiling | __shared__ float As[BM][BK]; | 3-5x |
| L1_PRESSURE_BOUND | Padding | As[BM][BK+1] | 1.2-2x |
| LATENCY_BOUND | Double Buffer | As[2][BM*BK] | 1.2-1.5x |
| COMPUTE_BOUND | FMA | fmaf(a, b, c) | 1.1-1.3x |
| OCCUPANCY_BOUND | 调整 block size | __launch_bounds__(256, 2) | 1.2-2x |
# 完整采集 (推荐)
ncu --set full -o report_name --target-processes all ./kernel
# 指定 sections
ncu --section SpeedOfLight,Occupancy,LaunchStats -o report_name ./kernel
# 特定指标
ncu --metrics sm__throughput.avg.pct,dram__throughput.avg.pct -o report_name ./kernel
# 查看摘要
ncu --import report.ncu-rep --print-summary per-kernel
# 查看详情
ncu --import report.ncu-rep --page details
# 导出 CSV
ncu --import report.ncu-rep --page raw --csv > metrics.csv
# 对比两个报告
ncu --diff report1.ncu-rep report2.ncu-rep
高 Throughput ≠ 高效率
DRAM Throughput 低可能是好事
Occupancy 不是越高越好
examples/本 Skill 支持完整的自动化 NCU 性能分析工作流，包含全量采集和持久化存储
Optimize Mixture-of-Experts (MoE) CUDA kernels on NVIDIA B200 (Blackwell, SM100). Use when writing, migrating, debugging, or optimizing FP8/FP16 MoE operators that involve grouped GEMM, per-expert routing, SwiGLU activation, gather/scatter, or Blackwell-specific features (TMA, tcgen05, TMEM). Covers PyTorch→CUDA migration, an optimization ladder ordered by measured ROI, FP8 correctness failure-mode catalog, cuBLAS/CUTLASS/tcgen05 backend selection, and agent-team patterns for multi-round optimization.
Multi-agent parallel CUDA kernel optimization. Use this skill when the user wants to run two (or more) optimization agents in parallel — typically on separate GPUs — to explore different optimization directions concurrently, or when the single-agent loop has plateaued and the user asks for parallel exploration. Also use when the user says things like "spawn agents on GPUs 6 and 7", "have two workers try different directions", or "Agent 1 focus on X, Agent 2 focus on Y". Extends `cuda-kernel-autodev` (which defines the single-agent loop) with orchestration protocol: isolated working directories, per-agent GPU pinning, sync/merge cadence, direction assignment, and main-agent role. Typically activated after you've already been running `cuda-kernel-autodev` for a while and the curve has flattened.
CUDA kernel development and optimization for NVIDIA B200 (Blackwell, compute capability 10.0) in Claude Code. Use when writing, reviewing, debugging, profiling, or optimizing CUDA kernels for B200/Blackwell. Covers kernel structure, memory hierarchy, shared memory and cluster constraints, PTX/TensorCore inspection, nsys/ncu profiling, compute-sanitizer debugging, and use of local PTX/Runtime/Driver reference docs. Prefer this skill for CUDA, PTX, Tensor Core, WGMMA/TMA/tcgen05, shared memory, occupancy, coalescing, register pressure, launch configuration, and Blackwell-specific tuning.
Autonomous end-to-end CUDA kernel development and optimization workflow. Use this skill whenever the user wants to write, optimize, or speed up a CUDA kernel/operator against a reference — whether they say "optimize this kernel", "improve GPU throughput", "beat PyTorch baseline", "speed up my CUDA op", "auto-tune this kernel", are iterating on kernel.cu / kernel.cpp / kernel.py in a benchmark framework (FlashInfer, KernelBench, MLSys contests), or are asking how to structure kernel optimization experiments. Covers the full 3-phase workflow: setup/baseline → hypothesize-implement-eval-commit loop with keep/revert discipline → submission. Explicitly pairs with `cuda-roofline-strategy` for choosing what to try next, `cuda-kernel-techniques` for the catalog of proven techniques, `cuda-agent-team` for multi-agent parallel mode, and `ncu-cuda-profiling` for collecting NCU data.
Catalog of proven CUDA kernel optimization techniques, broken by sub-topic (memory access, data placement, parallelism, compute, control flow, occupancy, numerical, anti-patterns). Use this skill whenever you need to look up *how* a specific technique works or *when* it applies — e.g., "how does Split-K FlashDecoding work", "when should I use cp.async vs TMA", "how do I implement online softmax", "what is register-centric data path", "what bank-conflict padding should I use", "when does warp specialization help", or before committing to a technique suggested by `cuda-roofline-strategy`. The main SKILL.md is an index; each sub-topic reference has the implementation details, when-it-helps, when-it-hurts, and code sketches. Each technique includes hardware-general guidance with B200 (SM100) and older-arch examples where they differ.
Roofline-driven strategy selection for CUDA kernel optimization. Use this skill during the "what do I try next" moment in a kernel optimization loop — when the user has a profile (NCU report or similar) and needs to decide which class of technique to try, or when experiments are plateauing and the user asks "what now?", or when diagnosing whether a kernel is compute-bound / bandwidth-bound / occupancy-limited / latency-limited. Maps (roofline position, iteration phase) → technique tier so the next experiment has a principled justification rather than being a guess. Complements `cuda-kernel-autodev` (which drives the loop) and `cuda-kernel-techniques` (which details each technique).