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

files

ships runnable scripts

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_* |

| 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/

Other skills for the same job

different authors, same section of the catalogue

Use Dom

by openai

vendor ×3

Use Expo DOM components to run web code in a webview on native and as-is on web. Migrate web code to native incrementally.

3k tokens

Frontend Dev Guidelines

by mrgoonie

×2

Frontend development guidelines for React/TypeScript applications. Modern patterns including Suspense, lazy loading, useSuspenseQuery, file organization with features directory, MUI v7 styling, TanStack Router, performance optimization, and TypeScript best practices. Use when creating components, pages, features, fetching data, styling, routing, or working with frontend code.

33k tokens

Next Cache Components

by vercel-labs

vendor ×2

Next.js 16 Cache Components - PPR, use cache directive, cacheLife, cacheTag, updateTag

| Build Shopify apps, extensions, themes using GraphQL Admin API, Shopify CLI, Polaris UI, and Liquid. "shopify theme", "liquid template", "polaris", "shopify graphql", "shopify webhook", "shopify billing", "app subscription", "metafields", "shopify functions"

Build Gradio web UIs and demos in Python. Use when creating or editing Gradio apps, components, event listeners, layouts, or chatbots.

10k tokens

Figma Generate Diagram

by openai

vendor ×1

MANDATORY prerequisite — load this skill BEFORE every `generate_diagram` tool call. NEVER call `generate_diagram` directly without loading this skill first. Trigger whenever the user asks to create, generate, draw, render, sketch, or build a diagram — flowchart, architecture diagram, sequence diagram, ERD or entity-relationship diagram, state diagram or state machine, gantt chart, or timeline. Also trigger when the user mentions Mermaid syntax or wants a system architecture, decision tree, dependency graph, API call flow, auth handshake, schema, or pipeline visualized in FigJam. Routes to type-specific guidance, sets universal Mermaid constraints, and tells you when to use a different diagram type or skip the tool entirely (mindmaps, pie charts, class diagrams, etc.).

Analyzes web performance using Chrome DevTools MCP. Measures Core Web Vitals (LCP, INP, CLS) and supplementary metrics (FCP, TBT, Speed Index), identifies render-blocking resources, network dependency chains, layout shifts, caching issues, and accessibility gaps. Use when asked to audit, profile, debug, or optimize page load performance, Lighthouse scores, or site speed. Biases towards retrieval from current documentation over pre-trained knowledge.

Algorithmic philosophies are computational aesthetic movements that are then expressed through code. Output .md files (philosophy), .html files (interactive viewer), and .js files (generative algorithms).

15k tokens scripts

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.