Loading...
Loading...
Write, debug, and optimize CUTLASS and CuTeDSL GPU kernels using local source code, examples, and header references. Use when the user mentions CUTLASS, CuTe, CuTeDSL, cute::Layout, cute::Tensor, TiledMMA, TiledCopy, CollectiveMainloop, CollectiveEpilogue, GEMM kernel, grouped GEMM, sparse GEMM, flash attention CUTLASS, blackwell GEMM, hopper GEMM, FP8 GEMM, blockwise scaling, MoE GEMM, StreamK, warp specialization CUTLASS, TMA CUTLASS, or asks about writing high-performance CUDA kernels with CUTLASS/CuTe templates.
npx skill4agent add slowlyc/agent-gpu-skills cutlass-skillrepos/cutlass/~/.cursor/skills/cutlass-skill/repos/cutlass/~/.claude/skills/cutlass-skill/repos/cutlass/~/.agents/skills/cutlass-skill/repos/cutlass/~/.cursor/skills/cutlass-skill/repos/cutlass/bash update-repos.sh cutlassCUTLASS_REPO/python/CuTeDSL/
├── cutlass/
│ ├── base_dsl/ # DSL 基础: 类型, 变量, 函数, PTX emit
│ ├── cute/ # CuTe Python 绑定: Layout, Tensor, TiledMMA, TiledCopy
│ ├── cutlass_dsl/ # CUTLASS DSL: GEMM builder, epilogue, pipeline
│ ├── pipeline/ # 流水线抽象: MainloopPipeline, PipelineAsync
│ ├── jax/ # JAX 集成
│ ├── utils/ # 编译工具, profiler, tensor 工具
│ └── torch.py # PyTorch 集成CUTLASS_REPO/examples/python/CuTeDSL/
├── ampere/ # Ampere: sgemm, tensorop_gemm, flash_attention_v2
├── hopper/ # Hopper: TMA gemm, FP8, grouped GEMM
├── blackwell/ # Blackwell: blockwise_gemm
├── blackwell_geforce/ # GeForce Blackwell
├── cute/ # CuTe tutorials (Python)
├── distributed/ # 分布式 GEMM
├── experimental/ # 实验性功能
├── jax/ # JAX 示例
├── notebooks/ # Jupyter notebooks (英文)
├── notebooks-zh/ # Jupyter notebooks (中文)
└── advanced_compiler_control/ # 高级编译控制CUTLASS_REPO/examples/
├── 00-47: Ampere 及更早架构
├── 48-69: Hopper (sm_90)
│ ├── 48_hopper_warp_specialized_gemm
│ ├── 49_hopper_gemm_with_collective_builder
│ ├── 54_hopper_fp8_warp_specialized_gemm
│ ├── 55_hopper_mixed_dtype_gemm
│ ├── 57_hopper_grouped_gemm
│ ├── 62_hopper_sparse_gemm
│ ├── 67_hopper_fp8..._blockwise_scaling
│ ├── 88_hopper_fmha
│ └── ...
├── 70-93: Blackwell (sm_100)
│ ├── 70_blackwell_gemm
│ ├── 71_blackwell_gemm_with_collective_builder
│ ├── 72_blackwell_narrow_precision_gemm
│ ├── 77_blackwell_fmha
│ ├── 81_blackwell_gemm_blockwise
│ ├── 83_blackwell_sparse_gemm
│ ├── 92_blackwell_moe_gemm
│ ├── 93_blackwell_low_latency_gqa
│ └── ...
└── cute/tutorial/ # CuTe C++ tutorials (sgemm, tiled_copy, hopper, blackwell)CUTLASS_REPO/include/cute/
├── layout.hpp # Layout 核心: Shape, Stride, 组合
├── tensor.hpp # Tensor: make_tensor, local_tile, partition
├── swizzle.hpp # Swizzle 模式
├── algorithm/ # copy, gemm, fill, clear
├── arch/ # 架构特定: copy_sm90, mma_sm90, copy_sm100
├── atom/ # MMA atom, Copy atom 定义
│ ├── mma_atom.hpp
│ ├── copy_atom.hpp
│ └── mma_traits_sm90_gmma.hpp # WGMMA traits
├── numeric/ # 数值类型
└── container/ # tuple, arrayCUTLASS_REPO/include/cutlass/
├── gemm/ # GEMM 设备层, collective, kernel
│ ├── collective/ # CollectiveMainloop, CollectiveEpilogue
│ ├── kernel/ # GemmUniversal
│ └── device/ # 设备启动接口
├── epilogue/ # Epilogue: bias, activation, scaling
├── conv/ # 卷积
├── arch/ # MMA 指令包装 (mma_sm90.h, mma_sm100.h)
├── pipeline/ # Pipeline: PipelineTmaAsync, PipelineAsync
├── experimental/ # 实验性 API
└── detail/ # 内部实现细节CUTLASS_REPO/python/pycute/
├── layout.py # Layout, make_layout, complement, coalesce
├── int_tuple.py # IntTuple 操作
├── swizzle.py # Swizzle
└── typing.py # 类型定义CUTLASS_REPO="$HOME/.cursor/skills/cutlass-skill/repos/cutlass"
# 查找 CuTeDSL GEMM 示例
rg "@jit" $CUTLASS_REPO/examples/python/CuTeDSL/
# 查找 TiledMMA 使用
rg "TiledMMA\|tiled_mma" $CUTLASS_REPO/python/CuTeDSL/cutlass/cute/
# 查找 pipeline 用法
rg "MainloopPipeline\|PipelineAsync" $CUTLASS_REPO/python/CuTeDSL/cutlass/pipeline/
# 查找 Blackwell CuTeDSL 示例
rg "sm_100\|blackwell" $CUTLASS_REPO/examples/python/CuTeDSL/blackwell/# 查找 Layout 操作
rg "make_layout\|composition\|complement" $CUTLASS_REPO/include/cute/layout.hpp
# 查找 TiledCopy 使用
rg "TiledCopy\|make_tiled_copy" $CUTLASS_REPO/include/cute/
# 查找 MMA atom traits
rg "MMA_Traits" $CUTLASS_REPO/include/cute/atom/
# 查找 Hopper WGMMA
rg "SM90_64x" $CUTLASS_REPO/include/cute/atom/mma_traits_sm90_gmma.hpp
# 查找 TMA copy
rg "SM90_TMA" $CUTLASS_REPO/include/cute/arch/# 查找 CollectiveBuilder 使用
rg "CollectiveBuilder" $CUTLASS_REPO/examples/49_hopper_gemm_with_collective_builder/
# 查找 Collective Mainloop
rg "CollectiveMainloop" $CUTLASS_REPO/include/cutlass/gemm/collective/
# 查找 Epilogue 融合
rg "fusion\|EVT" $CUTLASS_REPO/include/cutlass/epilogue/
# 查找 kernel 启动模板
rg "GemmUniversal" $CUTLASS_REPO/include/cutlass/gemm/device/# 查找 FP8 GEMM 配置
rg "float_e4m3\|float_e5m2\|fp8" $CUTLASS_REPO/examples/54_hopper_fp8_warp_specialized_gemm/
# 查找 blockwise scaling
rg "blockwise\|block_scale" $CUTLASS_REPO/examples/67_hopper_fp8_warp_specialized_gemm_with_blockwise_scaling/
# 查找 grouped GEMM
rg "grouped\|ProblemShape::Group" $CUTLASS_REPO/examples/57_hopper_grouped_gemm/
# 查找 sparse GEMM
rg "sparse\|Sparse" $CUTLASS_REPO/examples/62_hopper_sparse_gemm/
# 查找 StreamK
rg "StreamK\|stream_k" $CUTLASS_REPO/examples/47_ampere_gemm_universal_streamk/| Need | Source | Path |
|---|---|---|
| CuTeDSL 入门 | CuTeDSL examples | |
| CuTeDSL Hopper GEMM | CuTeDSL examples | |
| CuTeDSL Blackwell GEMM | CuTeDSL examples | |
| CuTeDSL API 定义 | CuTeDSL source | |
| CuTe Layout 语义 | CuTe headers | |
| CuTe Tensor 操作 | CuTe headers | |
| MMA atom traits | CuTe atom | |
| TMA copy 架构 | CuTe arch | |
| CUTLASS GEMM 模板 | CUTLASS examples | |
| Collective Builder | CUTLASS examples | |
| Collective Mainloop | CUTLASS headers | |
| Epilogue 融合 | CUTLASS headers | |
| Pipeline 抽象 | CUTLASS headers | |
| pycute Layout 操作 | pycute | |
| Hopper FMHA | CUTLASS examples | |
| Blackwell FMHA | CUTLASS examples | |
| MoE GEMM | CUTLASS examples | |
from cutlass import jit, Int32, Float32
@jit
def add_kernel(x: Float32, y: Float32) -> Float32:
return x + yexamples/python/CuTeDSL/ampere/elementwise_add.pyexamples/python/CuTeDSL/ampere/sgemm.pyexamples/python/CuTeDSL/hopper/examples/cute/tutorial/sgemm_1.cusgemm_sm80.cu# 编译 CUTLASS example
cd CUTLASS_REPO && mkdir -p build && cd build
cmake .. -DCUTLASS_NVCC_ARCHS=90a # Hopper
cmake --build . --target 49_hopper_gemm_with_collective_builder
# 编译特定架构
cmake .. -DCUTLASS_NVCC_ARCHS="80;90a;100a"
# CuTeDSL 运行
pip install -e python/CuTeDSL/
python examples/python/CuTeDSL/ampere/sgemm.py| 问题 | 可能原因 | 查找参考 |
|---|---|---|
| GEMM 精度不对 | Epilogue 未配置正确的 accumulator 类型 | |
| TMA 报错 | Tensor alignment 不满足 128B | |
| Collective Builder 编译失败 | 架构不匹配或 pipeline 配置错误 | |
| CuTeDSL 类型错误 | DSL 类型与 CUDA 类型不匹配 | |
| Layout swizzle 错误 | Swizzle 模式与数据排布不兼容 | |
# 在 cursor-gpu-skills 项目目录下
bash update-repos.sh cutlassCUTLASS_REPO/media/docs/cute/examples/python/CuTeDSL/notebooks/