kernel-skills

GPU kernel 开发相关的 Claude Code / Codex skill 集合。按编程模型分目录组织，每个子目录是一个可独立使用的 skill。

目录结构

kernel-skills/
├── cuda/                            # NVIDIA CUDA 相关 skill
│   ├── cute-dev-setup/              # 配置 CUTLASS/CuTe 开发环境
│   ├── operator-accuracy-test/      # 算子精度测试（数值对拍）标准流程
│   ├── operator-performance-test/   # 算子 Quick/Full 性能测试标准流程
│   ├── kernel-study/                # 交互式 CuTe kernel 教学（多轮课程）
│   ├── kernel-docgen/               # 把 kernel 优化记录为互链双文档
│   └── kernel-opt-loop/             # kernel 优化闭环与跨会话决策记忆
├── triton/                          # Triton 相关 skill（占位）
└── npu/                             # 国产 NPU（昇腾等）相关 skill（占位）

怎么串起来用

六个 skill 不是并列的工具箱，而是一条链：环境 → 读懂 → 证据 → 优化 → 沉淀。

flowchart LR
    S["cute-dev-setup<br/>环境就位"] --> T["kernel-study<br/>读懂 kernel"]
    T --> L["kernel-opt-loop<br/>优化闭环"]
    L --> D["kernel-docgen<br/>正式文档"]
    A["operator-accuracy-test<br/>精度 gate"] --> L
    P["operator-performance-test<br/>性能证据"] --> L
    S --> A
    S --> P
    T -. 结课 .-> D

Loading

• cute-dev-setup 是一次性的起点。编不过、IDE 报红的时候先解决它，否则后面每一步都会被环境问题干扰。
• kernel-study 用在接手一个看不懂的 kernel 时。它不产出代码改动，只把 kernel 讲明白，结课时留下一份闭卷试卷。
• 两个测试 skill 是优化的前置条件，不是收尾工作。精度测试提供正确性 gate，性能测试提供命名基线和 MFU/MBU/Roofline 数据。kernel-opt-loop 明确要求「无 NCU/benchmark 证据不立项」——没有这两样，优化就只能靠感觉。
• kernel-opt-loop 是优化的主循环，跨会话记住哪些方向已经试过、为什么被舍弃。它的记忆文件只记决策和数据，不是交付物。
• kernel-docgen 是唯一的正式出口。kernel-study 结课后和 kernel-opt-loop 跑完之后，都由它产出对外的中文文档。

只想解决眼前一件事：环境配不好用 cute-dev-setup，写完算子不放心用 operator-accuracy-test，想知道快慢用 operator-performance-test。这三个都可以独立使用，不必先跑完整条链。

已有 skill

cuda/cute-dev-setup

在标准 VSCode 环境下配置 CUTLASS/CuTe 的 CUDA 开发环境。CUTLASS 是 header-only，这个 skill 解决的是「怎么把它接进你的项目和 IDE」。

包含：

• 依赖检测脚本 — 自动探测 nvcc 路径（避开 /usr/bin/nvcc 的 bug）、GCC 完整性（检查 cc1plus 是否残缺）、GPU 架构、CUTLASS 是否就位、libcu++ 头文件
• 编译脚本模板 — 封装 nvcc 调用，按 GPU compute capability 选 sm_xx，正确处理 -ccbin/--allow-unsupported-compiler/-extended-lambda
• VSCode IntelliSense 配置模板 — 解决 cuda/std/utility、stddef.h 找不到、__global__ 报红等常见问题
• 故障排查手册 — 按 nvcc/cc1plus/VSCode 报错关键词查表的诊断指南
• 最小验证程序 — 确认环境能编译能跑的 hello world

cuda/operator-accuracy-test

为手写 CUDA/CuTe 算子落地精度测试标准流程（数值对拍）。解决三个痛点：kernel 复制粘贴进测试导致源改测试过期假绿、不知 reference 用什么/容差定多少、gemm 用 cuBLAS 当 reference 被 TF32 静默污染。

包含：

• SOP 模板 — 12 节标准流程（reference 选型、多 dtype 容差表、TF32 污染规避、NaN 处理、退出码语义、报告格式、kernel 同步机制、扩展点），复制到项目根即用
• 测试脚手架脚本 — ./new_accuracy_test.sh <op> 生成 test_<op>.cu 骨架，模板化 compareArrays<T>（支持 fp32/bf16/fp16/fp8）+ tolerance_for 容差查表 + 对齐 PyTorch 的报告格式 + 退出码语义
• .cuh 范式样板 — 对齐 CUTLASS 官方的 kernel 头文件范式（#pragma once + namespace + 函数内 using namespace），通过 #include 机制杜绝复制粘贴分叉
• vecAdd 端到端范例 — 可直接编译运行的三 case 测试（含 .cuh + test.cu），照着改新算子
• 容差与 TF32 事实依据 — 每个容差值和 TF32 处理附 PyTorch/cuBLAS/nvcc 官方文档出处

第一阶段做数值误差对拍（稳定性/跨硬件留扩展点）；目标 GPU / 架构由项目 SOP 声明。

cuda/operator-performance-test

为 CUDA/CuTe 算子建立可持续的 Quick/Full 性能测试标准流程。Quick 使用 CUDA Event 提供延迟统计、MFU/MBU、analytical Roofline、reference speedup 和命名基线回归判断；Full 运行固定分层 case，并用 NSYS/NCU 深入分析一个预先配置的代表 case。

包含：

• Quick/Full 执行契约 — 正确性 gate、warmup/自动 batching、median/P10/P90/min、固定 case 分层与代表 case profiling
• 指标语义 — 区分 MFU/MBU/analytical Roofline 派生估算和 NCU 硬件计数器实测结论
• 命名基线与回归门禁 — 显式 create/update、Git 可追踪，默认 >5% 且 >0.2 us，连续两次命中才判定回归
• 安全的 NCU 权限流程 — Full 只交互认证一次，不保存密码，结束后清除 sudo ticket
• 状态与报告规范 — PASS、REGRESSION、UNSTABLE、PROFILE_PARTIAL、INVALID，以及 JSON/Markdown/manifest 产物约定
• SOP 与算子配置模板 — 可复制到项目中扩展新算子

目标 GPU / 架构由项目 SOP 声明；换卡或换架构时需要显式扩展和重新校准。

cuda/kernel-study

把一个 CuTe kernel 讲成一门多轮互动课程：先校准学习者背景（expert / cuda-only / beginner）和讲法（top-down / skeleton-first / by-technique），L0 概览后每轮只讲一课，停下等「继续」。教学重点是优化思维（每个优化点回答 Where/What/Why/Counterfactual），而非逐行复述源码。

包含：

• 完整课程工作流 — 读真实源码 → 校准 → L0 概览（数据流图 + actor 模型 + 优化点清单 + 大纲）→ 每轮一课 → 结课总表
• 产物生命周期管理 — 大纲与每课记录持久化到 docs/kernel-study/<kernel>/（工作目录，结课删除）；结课时合成闭卷试卷 + 答案（docs/kernel-notes/<kernel>_exam.md / _answer.md，永久交付）
• kernel 家族特例 — 共享大量机制的姊妹 kernel（同一算子的若干量化/调度变体）推迟到家族末尾出对比试卷，避免重复考核
• Lesson menu + 概念库 — 按代码区域（数据格式 / config / 骨架 / producer / consumer / mainloop / epilogue / 调度 / launcher）组织，CuTe 概念（TiledMMA、TMA、mbarrier、warp specialization 等）即用即讲
• 结课后 handoff 到 kernel-docgen 生成正式优化文档

cuda/kernel-docgen

把一次 kernel 优化记录为两篇互链的中文活文档：算子优化文档（实例层，docs/kernel-opt/<kernel>.md）+ 技术优化文档（方法/类别层，docs/kernel-notes/<technique>.md），形成 方法 → 类别 → 实例 三层知识体系。

包含：

• 双产物工作流 — 7 步流程：解析 slug → 收集指标 → 归类技术/类别 → 填算子文档 → 填/更新技术文档 → 互链 → 汇报
• 两份成稿模板 — 算子文档 + 技术文档待填模板，即两类文档的权威结构源
• 伪代码核心规则 — 实现思路用语法合法的伪代码骨架表达思想，禁止逐字粘贴 kernel 源码
• 指标自动采集 — 从工作区的优化迭代记录 / bench 输出 / git log 提取优化叙事与数值，缺定量字段才问用户
• Living-doc 更新规则 — 同算子/同技术复用原文档就地更新，追加历史而非重建

cuda/kernel-opt-loop

规范 CUDA kernel 优化的跨会话闭环：假设 → 实现 → 正确性与 paired benchmark 验证 → 保留/舍弃/搁置决策 → 更新长期记忆。解决跨会话重复踩坑、凭对话印象报性能数字、优化上下文随会话丢失的问题。

包含：

• 五段记忆文件模板 — 当前状态快照、数据契约与上下文、瓶颈清单、append-only 决策日志、已舍弃/搁置方向（含「重新考虑的条件」）
• 阶段化工作流 — 接手先读记忆汇报快照；改代码前先立项假设（无 NCU/benchmark 证据不立项，默认收益门槛几何平均 ≥3%）；一次只推进一个假设
• 验证与决策规范 — 正确性前置、paired ABBA 交替对比、环境信息（GPU/CUDA/驱动/commit）缺一不可、三态决策必须附量化数据
• 防重复踩坑机制 — 提议命中已舍弃/搁置方向时先出示当时的原因，重新考虑的条件不满足不建议重启

记忆文件只记录决策与数据，不替代交付文档；正式优化文档仍由 kernel-docgen 产出。

使用方式

每个 skill 就是一个带 frontmatter 的 SKILL.md 目录。Claude Code 和 Codex 共用这套格式，同一份目录拷到哪边都能用，区别只在目标路径。

装到 Claude Code：

git clone https://github.com/yiwen-cai/kernel-skills.git
mkdir -p ~/.claude/skills
cp -r kernel-skills/cuda/* ~/.claude/skills/

Claude Code 会自动发现。

装到 Codex：

git clone https://github.com/yiwen-cai/kernel-skills.git
mkdir -p "${CODEX_HOME:-$HOME/.codex}/skills"
cp -r kernel-skills/cuda/* "${CODEX_HOME:-$HOME/.codex}/skills/"

Codex 的 skill 目录由 $CODEX_HOME 决定，未设置时默认 ~/.codex。装完要重启 Codex 才会加载。

上面两段都是全量安装六个 skill。只要其中几个就把 cuda/* 换成具体目录名：

cp -r kernel-skills/cuda/cute-dev-setup ~/.claude/skills/
cp -r kernel-skills/cuda/operator-accuracy-test ~/.claude/skills/

装好之后说「配置 cute 环境」「.cu 文件报红」「给算子写精度测试」等就会自动触发。

手动用脚本/模板：skill 目录里的 assets/ 可以脱离 agent 单独使用，比如直接拷 new_accuracy_test.sh 生成测试骨架、拷 SOP.md.tmpl 落地项目 SOP。

相关项目与技术分享

kernel-examples — 手撕算子模板库，面试高频题的参考实现。本仓库给的是流程（怎么配环境、怎么测、怎么优化、怎么沉淀文档），kernel-examples 给的是具体代码。

按 栈 → 算子 → 文件 三层组织，同一个算子可以跨语言对照：

• CUDA — LayerNorm（two-pass 基线与 Welford single-pass 两版对照）、RMSNorm、float4 向量化转置（shared memory 消 bank conflict）、inclusive/exclusive scan、int8 per-token 量化
• CuTe — SM80 BF16 TN GEMM
• Triton — LayerNorm、RMSNorm、softmax、FlashAttention forward、scan、int8 per-token / per-group 量化

它自带一套 Mac 友好的 LSP 配置（bootstrap_lsp.sh 把 CUDA / CUTLASS / Triton 头文件拉到本地），但只保证跳转和补全，不能编译也不能跑。想真正编译、对精度、测性能，正是本仓库 cute-dev-setup 和两个测试 skill 接手的地方。

yiwen-cai.github.io — 我的技术博客，持续更新，写的范围比这两个仓库宽：算子只是其中一块，还有 LLM 推理、分布式训练和论文精读。分类入口：CUDA · Triton · LLM 推理 · 系统 · 论文。

几篇代表性的长文：

• 长上下文 LLM 推理中的 KV Cache 优化综述 — 系统管理、缓存压缩与架构协同
• RL Meets LLMs — 大语言模型全生命周期强化学习综述
• Kernels、Triton 与 Profiling — CS336 Lecture 6 笔记，和本仓库的 profiling 流程对得上

模型 Provider

跑这些 skill 需要自备模型 API。推荐 OrcaRouter：一个 OpenAI 兼容的 AI 网关，单个 key 接 200+ 模型，token 按上游厂商公开价结算不加价，带自动 failover，官方支持 Claude Code 和 Codex。

Claude Code 接入（base URL 不带 /v1，Anthropic 协议的客户端会自己补 /v1/messages）：

export ANTHROPIC_BASE_URL=https://api.orcarouter.ai
export ANTHROPIC_AUTH_TOKEN=sk-orca-your-key
export ANTHROPIC_API_KEY=""

模型 ID 用 vendor/model 格式（如 anthropic/claude-opus-4.8），或者用 orcarouter/auto 让它按 prompt 自动选。

License

MIT