mcpbeat

Cutlass Skill

slowlyc/cutlass-skill

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, FP4 GEMM, blockwise scaling, MoE GEMM, StreamK, warp specialization CUTLASS, TMA CUTLASS, epilogue fusion, EVT (Epilogue Visitor Tree), pycute, Layout algebra, Swizzle pattern, GemmUniversal, KernelSchedule, EpilogueSchedule, CUTLASS collective builder, CUTLASS pipeline, or asks about writing high-performance CUDA kernels with CUTLASS/CuTe templates. Also use when the user wants to understand CUTLASS source code structure, compile CUTLASS examples, or debug CUTLASS template errors.

3k tokens
context cost
the whole folder, loaded on every use
2
files
ships runnable scripts
0
copies elsewhere
how many repositories repackaged it
153
stars on the repo
on the repository, not the skill itself

Install

one command, takes just this skill from the repository 
npx skills add https://github.com/slowlyC/agent-gpu-skills --skill cutlass-skill

The instruction itself

20 sections, as written by the author

CUTLASS & CuTeDSL Development

Source Code Locations

CUTLASS 源码位于此 skill 安装目录下的 repos/cutlass/。

实际路径取决于所用工具:

  • Cursor: ~/.cursor/skills/cutlass-skill/repos/cutlass/
  • Claude Code: ~/.claude/skills/cutlass-skill/repos/cutlass/
  • Codex: ~/.codex/skills/cutlass-skill/repos/cutlass/

CUTLASS_REPO: 下文示例用 ~/.cursor/skills/cutlass-skill/repos/cutlass/ 作占位符，替换为实际路径。

如果该路径不存在，在项目目录下运行 bash update-repos.sh cutlass。

CuTeDSL (Python DSL for GPU Kernels)

CUTLASS_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 集成

CuTeDSL 示例:

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 C++ Examples (按架构分类)

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)

CuTe C++ Headers

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, array

CUTLASS C++ Headers

CUTLASS_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/             # 内部实现细节

pycute (Python CuTe 绑定)

CUTLASS_REPO/python/pycute/
├── layout.py           # Layout, make_layout, complement, coalesce
├── int_tuple.py        # IntTuple 操作
├── swizzle.py          # Swizzle
└── typing.py           # 类型定义

Search Strategy

用 Grep 工具搜索，不要整文件加载。

CuTeDSL 用法

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/

CuTe C++ 用法

# 查找 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/

CUTLASS Collective Builder

# 查找 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/

GEMM 示例搜索

# 查找 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/

When to Use Each Source

| Need | Source | Path |

|------|--------|------|

| CuTeDSL 入门 | CuTeDSL examples | examples/python/CuTeDSL/ampere/ |

| CuTeDSL Hopper GEMM | CuTeDSL examples | examples/python/CuTeDSL/hopper/ |

| CuTeDSL Blackwell GEMM | CuTeDSL examples | examples/python/CuTeDSL/blackwell/ |

| CuTeDSL API 定义 | CuTeDSL source | python/CuTeDSL/cutlass/ |

| CuTe Layout 语义 | CuTe headers | include/cute/layout.hpp |

| CuTe Tensor 操作 | CuTe headers | include/cute/tensor.hpp |

| MMA atom traits | CuTe atom | include/cute/atom/ |

| TMA copy 架构 | CuTe arch | include/cute/arch/copy_sm90* |

| CUTLASS GEMM 模板 | CUTLASS examples | examples/48-93_* |

| Collective Builder | CUTLASS examples | examples/49_hopper_gemm_with_collective_builder/ |

| Collective Mainloop | CUTLASS headers | include/cutlass/gemm/collective/ |

| Epilogue 融合 | CUTLASS headers | include/cutlass/epilogue/ |

| Pipeline 抽象 | CUTLASS headers | include/cutlass/pipeline/ |

| pycute Layout 操作 | pycute | python/pycute/ |

| Hopper FMHA | CUTLASS examples | examples/88_hopper_fmha/ |

| Blackwell FMHA | CUTLASS examples | examples/77_blackwell_fmha/ |

| MoE GEMM | CUTLASS examples | examples/92_blackwell_moe_gemm/ |

CuTeDSL 编写模式

基本 Elementwise Kernel

from cutlass import jit, Int32, Float32

@jit
def add_kernel(x: Float32, y: Float32) -> Float32:
    return x + y

参考 examples/python/CuTeDSL/ampere/elementwise_add.py 获取完整示例。

CuTeDSL GEMM

参考 examples/python/CuTeDSL/ampere/sgemm.py 获取基础 SGEMM。

参考 examples/python/CuTeDSL/hopper/ 获取 Hopper TMA GEMM。

CuTe C++ GEMM 模式

参考 examples/cute/tutorial/sgemm_1.cu ~ sgemm_sm80.cu 获取 CuTe SGEMM 渐进教程。

Compilation Reference

# 编译 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 类型 | rg "ElementAccumulator" examples/49_* |

| TMA 报错 | Tensor alignment 不满足 128B | rg "Alignment\|alignment" examples/48_* |

| Collective Builder 编译失败 | 架构不匹配或 pipeline 配置错误 | rg "KernelSchedule\|EpilogueSchedule" examples/49_* |

| CuTeDSL 类型错误 | DSL 类型与 CUDA 类型不匹配 | rg "dtype\|element_type" examples/python/CuTeDSL/ |

| Layout swizzle 错误 | Swizzle 模式与数据排布不兼容 | include/cute/swizzle.hpp |

更新 CUTLASS 源码

# 在 cursor-gpu-skills 项目目录下
bash update-repos.sh cutlass

Additional References

  • CUTLASS 官方文档: https://github.com/NVIDIA/cutlass
  • CuTe 文档: CUTLASS_REPO/media/docs/cute/ (如果使用 --full 模式安装)
  • CuTeDSL notebooks: examples/python/CuTeDSL/notebooks/

How to use it

Copy the folder

Take slowlyc/cutlass-skill from the repository into ~/.claude/skills for personal use, or into .claude/skills inside a project.

Check the name does not clash

The agent identifies a skill by the name field in its header. Two skills with the same name cannot sit side by side — one of them will be ignored.

Install what it needs

The instructions reference pip. Without those the skill loads but fails at the first command.