Cuda Cpp Kernel
vipshop/cache-dit
A skill your agent uses when writing, debugging, porting, reviewing, or optimizing CUDA C++ or PTX kernels; investigating CUDA Runtime or Driver API behavior; profiling kernels with Nsight Systems…
LeetCUDA 中文技术书（584 页，XeLaTeX 源）按需查阅 skill——写、优化、调试或 review CUDA C++/PTX kernel 时的权威参考路由层。当任务涉及：GPU 架构/Roofline/ occupancy、向量化与 coalescing、warp/block reduce、softmax（online/LSE merge）、…
Install Leetcuda Cpp Kernel
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a claude-codeProject install by default; add -g for ~/.claude/skills/.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernel --agent claude-codeProject scope by default; add --scope user for a personal install. Needs GitHub CLI 2.90.0 or later (public preview).
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .claude/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .claude/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.claude/skills/ instead of .claude/skills for a personal install. The folder must contain SKILL.md.
Claude Code skills documentation · loads skills from .claude/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .claude/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.Claude Code copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
$skill-installer install https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernelType this inside Codex. $skill-installer <name> installs a curated skill from openai/skills. The installer writes to $CODEX_HOME/skills (default ~/.codex/skills). Restart Codex if the skill does not show up.
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a codexProject install goes to .agents/skills/; add -g for ~/.codex/skills/.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernel --agent codexProject scope by default (.agents/skills/); add --scope user for a personal install.
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .agents/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .agents/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.agents/skills/ instead of .agents/skills for a personal install.
Codex skills documentation · loads skills from .agents/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .agents/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.Codex copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a cursorProject install goes to .agents/skills/; add -g for ~/.cursor/skills/.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernel --agent cursorProject scope by default (.agents/skills/); add --scope user for a personal install.
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .cursor/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .cursor/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.cursor/skills/ instead of .cursor/skills for a personal install.
Cursor skills documentation · loads skills from .cursor/skills/, .agents/skills/, .claude/skills/, .codex/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .cursor/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.Cursor copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
$ gemini skills install https://github.com/xlite-dev/LeetCUDA.git --path kernels/interview/book/skills/leetcuda-cpp-kernel--scope user (default) or --scope workspace; --path is the subfolder of the repo that holds the skill; --consent skips the security confirmation prompt.
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a gemini-cliProject install goes to .agents/skills/; add -g for ~/.gemini/skills/.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernel --agent gemini-cliProject scope by default (.agents/skills/); add --scope user for a personal install.
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .gemini/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .gemini/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.gemini/skills/ instead of .gemini/skills for a personal install, then run /skills reload.
Gemini CLI skills documentation · loads skills from .gemini/skills/, .agents/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .gemini/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.Gemini CLI copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernelInstalls for Copilot at project scope by default; add --scope user for a personal install. Preview a skill first with gh skill preview. Needs GitHub CLI 2.90.0 or later (public preview).
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a github-copilotProject install goes to .agents/skills/; add -g for ~/.copilot/skills/.
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .github/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .github/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.copilot/skills/ instead of .github/skills for a personal install. Commit .github/skills so cloud agent and code review can use it.
GitHub Copilot skills documentation · loads skills from .github/skills/, .claude/skills/, .agents/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .github/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.GitHub Copilot copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
$ npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a opencodeOpenCode documents no install command of its own. Project install goes to .agents/skills/; add -g for ~/.config/opencode/skills/.
$ gh skill install xlite-dev/LeetCUDA leetcuda-cpp-kernel --agent opencodeProject scope by default (.agents/skills/); add --scope user for a personal install.
$ git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git skills-src && mkdir -p .opencode/skills && cp -r skills-src/kernels/interview/book/skills/leetcuda-cpp-kernel .opencode/skills/leetcuda-cpp-kernel && rm -rf skills-srcUse ~/.config/opencode/skills/ instead of .opencode/skills for a personal install.
OpenCode skills documentation · loads skills from .opencode/skills/, .claude/skills/, .agents/skills/
Install the "leetcuda-cpp-kernel" agent skill from https://github.com/xlite-dev/LeetCUDA/tree/main/kernels/interview/book/skills/leetcuda-cpp-kernel into .opencode/skills/leetcuda-cpp-kernel/ in this project. Copy the whole folder (SKILL.md and every file beside it), keep the folder name "leetcuda-cpp-kernel", then confirm the skill loads.OpenCode copies the folder itself, the same result as the manual copy. Check what it changed before you commit it.
Facts
- Skill name
leetcuda-cpp-kernel- GitHub stars
- 12k
- Token cost
- ~3.5k tokens
- SKILL.md length
- 1,071 words
- Files
- 1
- Skills in repo
- 2
- Repo updated
- First seen
- Licence
- GPL-3.0
At a glance
LeetCUDA 中文技术书（584 页，XeLaTeX 源）按需查阅 skill——写、优化、调试或 review CUDA C++/PTX kernel 时的权威参考路由层。当任务涉及：GPU 架构/Roofline/ occupancy、向量化与 coalescing、warp/block reduce、softmax（online/LSE merge）、…
- Works in 4 steps: 环境变量 LEETCUDA_DIR（若已设置则直接使用） → 从当前工作目录及其父目录、兄弟目录中查找名为 LeetCUDA/ 且含 → 常见位置探测：~/LeetCUDA → …
- Tasks that involve GPU and accelerator computing
- SKILL.md covers 第一步：解析, 第二步：任务路由表, 第三步：按需读取纪律 and 已知重要结论速查（细节见对应章节）
- Calls pdftotext, git and pdftoppm; reaches github.com
What it does
Leetcuda Cpp Kernel is an agent skill from xlite-dev/LeetCUDA. LeetCUDA 中文技术书（584 页，XeLaTeX 源）按需查阅 skill——写、优化、调试或 review CUDA C++/PTX kernel 时的权威参考路由层。当任务涉及：GPU 架构/Roofline/ occupancy、向量化与 coalescing、warp/block reduce、softmax（online/LSE merge）、 SGEMV/SGEMM/HGEMM 阶梯优化、mma.sync/ldmatrix/WMMA、XOR/block swizzle、cp.async 多级流水、TMA/mbarrier/WGMMA（Hopper）、SM120 TMA+warp specialization、 FlashAttention FA1/FA2/FA3 实现、SM120 持久化 FA（persist-D 超越 cuDNN）、 split-D 大 headdim、CuTe Layout/Tensor/ TiledCopy/TiledMMA（含官方白皮书译注、colfax 范畴论译注、cute-zhihu 合集三份 深度参考）、FP8/FP4 量化 GEMM（e4m3…
Its SKILL.md is about 3.5k tokens, which your agent loads only when the skill is triggered. It is a single SKILL.md file with no bundled scripts.
It sits in AI & LLM Engineering, covering GPU and accelerator computing, CSS and styling and LaTeX. It works with CUDA, C++ and LaTeX. The repository describes itself as: An Open-source book with Modern CUDA learning Notes for Beginners - BF16/FP8/FP4, HGEMM, CuTe, Flash-Attention, etc. The licence is GPL-3.0.
When your agent uses it
- Tasks that involve GPU and accelerator computing
- Tasks that involve CSS and styling
- Tasks that involve LaTeX
Example prompts
- “任务 → 章节/图”
- “/leetcuda-cpp-kernel”
Workflow steps
4 steps, taken from the first numbered list in SKILL.md.
- 环境变量 LEETCUDA_DIR（若已设置则直接使用）
- 从当前工作目录及其父目录、兄弟目录中查找名为 LeetCUDA/ 且含
- 常见位置探测：~/LeetCUDA
- 本地没有则 clone（浅克隆即可）
What it can do on your machine
Read from SKILL.md and the folder at commit c5e985f. It shows what the files ask for, not the result of running them.
Tool permissions
Pre-approves nothing: there is no allowed-tools line, so your agent's usual permission prompts apply.
From allowed-tools in the SKILL.md frontmatter.
Runs code
Shell commands in SKILL.md call:
pdftotextgitpdftoppm
From the folder's file list and the shell code blocks in SKILL.md.
Network
Hosts in commands or code, which the agent is likely to contact:
github.com
From URLs in SKILL.md, links to its own repository left out.
Credentials
Names no API keys, tokens, secrets or passwords.
From names ending in _API_KEY, _TOKEN, _SECRET, _KEY or _PASSWORD in SKILL.md.
Context cost
Leetcuda Cpp Kernel loads about 3.5k tokens when it runs. Until then it costs about 225 tokens; SKILL.md has 1,071 words of instructions outside code blocks.
Estimates: characters ÷ 4, the usual rule of thumb; real counts depend on the model's tokenizer. Scripts and assets cost tokens only if the agent reads them.
Safety
Auto-check passed
The automated check found no risky patterns in SKILL.md.
Automated static check — not a guarantee. Review scripts before installing. It scans the text of SKILL.md for risky patterns (piping downloads into a shell, reading credential files, hidden Unicode, destructive commands); files beside SKILL.md are not scanned.
SKILL.md
The full file from xlite-dev/LeetCUDA at commit c5e985f, republished under its GPL-3.0 licence (© xlite-dev). 1,071 words, ~3,524 tokens.
Download SKILL.mdSave it as
.claude/skills/leetcuda-cpp-kernel/SKILL.md (or your agent's skills folder).- name
- leetcuda-cpp-kernel
- description
- LeetCUDA 中文技术书（584 页，XeLaTeX 源）按需查阅 skill——写、优化、调试或 review CUDA C++/PTX kernel 时的权威参考路由层。当任务涉及：GPU 架构/Roofline/ occupancy、向量化与 coalescing、warp/block reduce、softmax（online/LSE merge）、 SGEMV/SGEMM/HGEMM 阶梯优化、mma.sync/ldmatrix/WMMA、XOR/block swizzle、cp.async 多级流水、TMA/mbarrier/WGMMA（Hopper）、SM120 TMA+warp specialization、 FlashAttention FA1/FA2/FA3 实现、SM120 持久化 FA（persist-D 超越 cuDNN）、 split-D 大 head_dim、CuTe Layout/Tensor/ TiledCopy/TiledMMA（含官方白皮书译注、colfax 范畴论译注、cute-zhihu 合集三份 深度参考）、FP8/FP4 量化 GEMM（e4m3 per-row/per-block 前处理、NVFP4 两级量化、 block-scaled MMA m16n8k64.kind::mxf4nvf4 与 SF 流水）、FP8/FP4(NVFP4) 量化注意力、 nsys/ncu 性能分析、cuobjdump/ PTX/SASS 取证（含 setmaxnreg 与 shared::cluster/cta 陷阱）、CUDA 面试题时使用。 生产级 attention 算子参考首选 ffpa-attn（sm_120 CuTe fp16/fp8/fp4 峰值性能： fp16 超 FA-2、fp8 持平略优 Sage2、fp4 明显优于 Sage3，含 large head_dim 专门 优化）。本 skill 不复述书的内容，只做"任务 → 章节/图"的路由，agent 按需读取 书源文件。
- user-invocable
- true
leetcuda-cpp-kernel — LeetCUDA 书按需查阅（薄转发层）
本 skill 是 LeetCUDA 开源技术书 的路由层：书本身就是参考资源（chapters/*.tex 正文 + figures 配图 + appendices 附录 + colfax/cute-zhihu 两份独立深度参考文档）， skill 只负责把任务路由到正确的章节，不重写、不复述书内已有内容。
第一步：解析 <LeetCUDA_DIR>
按以下顺序定位 LeetCUDA 仓库根（记为 <LeetCUDA_DIR>），全程环境无关：
- 环境变量
LEETCUDA_DIR（若已设置则直接使用） - 从当前工作目录及其父目录、兄弟目录中查找名为
LeetCUDA/且含kernels/interview/的目录（如./LeetCUDA、../LeetCUDA、monorepo 内 兄弟目录布局） - 常见位置探测：
~/LeetCUDA - 本地没有则 clone（浅克隆即可）：
git clone --depth 1 https://github.com/xlite-dev/LeetCUDA.git <工作区>/LeetCUDA
书根（本 skill 全部资源相对它描述）：
<LeetCUDA_DIR>/kernels/interview/book目录结构：
| 路径（相对书根） | 内容 |
|---|---|
chapters/chNN-<slug>.tex | 41 个章节文件：ch00–ch37 主线 38 章 + ch19b 白皮书导读 + ch26b/ch26c 增补章 + _template.tex 章模板（中文，XeLaTeX 源；含行号锚定的源码解析、踩坑记录、实测性能表） |
chapters/wp/wp0-7.tex | ch19b 分节正文：CuTe 官方白皮书（Cris Cecka, arXiv:2603.02298）完整中文译注 |
figures/drawio/fig-<chNN>-<n>-<slug>/ | 历史留档：旧版 .drawio/gen.py/*.png 三件套（章内图已全部迁 TikZ inline，仅作旧版对照，勿当现行图源） |
figures/ffpa/、figures/tikz/、figures/misc/ | ffpa bench 图留档 / wp6 白皮书 2 张官方架构截图 / 封面 cover.pdf 等杂项 |
colfax/ | 独立译注文档：Colfax《Categorical Foundations for CuTe Layouts》（arXiv:2601.05972）中译，源在 colfax/sec/，成品 colfax/colfax-cute-zh.pdf |
cute-zhihu/ | 独立合集文档：reed（13 篇）+ 竹熙佳处（7 篇）CuTe 知乎系列忠实整理，源在 cute-zhihu/chapters/，成品 cute-zhihu/cute-zhihu.pdf |
references/ | zhihu-inventory.md（附录 E 数据源）+ fulltext/（7 篇知乎全文 markdown，可直接 grep） |
appendices/appA..appE | 见下文附录路由 |
tests/ | 每章最小正确性测试（CPU fp64 对拍，build_tests.sh --arch sm_120a --all） |
book.pdf | 编译成品（584 页，可直接 pdftotext 按页抽取） |
第二步：任务路由表
先按任务关键词找到章，再读该章 tex（先 grep -n '\\\\section' 看目录，再按行区间读，
不要整章全文读入）。全书章内配图已 100% 迁为 TikZ inline（2026-09-28 核实：
ch00–ch37 每章 2-11 个 tikzpicture，chapters/ 内 \includegraphics 全部为
% was: 历史注释；全书活跃图片引用仅 3 处——封面 figures/misc/cover.pdf +
ch19b/wp6 的 2 张官方架构截图），现行图源都在章节 tex 内。
Part I-II：CUDA 基础与 GEMM 阶梯（配图全部 TikZ 内联在各章 tex 内）
| 任务/症状 | 章 |
|---|---|
| 优化前方法论：先 nsys 后 ncu、torch.profiler、warmup/repeat 口径 | ch00-profiling.tex |
| 架构/执行模型/warp 调度/Roofline/Occupancy、寄存器与 255 墙 | ch01-arch-roofline.tex |
| warp shuffle 蝴蝶归约、block reduce、dot product | ch02-reduce-dot.tex |
| 向量化（f32x4/f16x8 pack）、coalescing、原子操作 | ch03-vectorize-atomic.tex |
| softmax 三级递进（naive→safe→online） | ch04-softmax.tex |
| LSE 与分块合并 merge_attn_states | ch05-lse-merge.tex |
| RMSNorm/LayerNorm（单遍统计） | ch06-norm.tex |
| RoPE、矩阵转置、bank conflict 与 padding | ch07-rope-transpose.tex |
| SGEMV 三种划分（memory-bound 专题） | ch08-sgemv.tex |
| SGEMM 阶梯一：block tile/Vec4/thread tile、双缓冲 | ch09-sgemm-tiling.tex |
| SGEMM 阶梯二：TF32 WMMA、cp.async 多级流水 | ch10-sgemm-tf32-wmma.tex |
| HGEMM：mma.sync m16n8k16、ldmatrix、fragment 布局 | ch11-hgemm-mma.tex |
| XOR swizzle、寄存器双缓冲、block swizzle、L2 复用 | ch12-hgemm-swizzle.tex |
Part III：Hopper/SM120 与 FlashAttention（配图全部 TikZ 内联）
| 任务/症状 | 章 |
|---|---|
| TMA（CUtensorMap/cp.async.bulk.tensor）、mbarrier 协议、WGMMA descriptor | ch13-hopper-tma-wgmma.tex |
| SM120：TMA+mma.sync+warp specialization、setmaxnreg 池数学与 shared::cluster/cta 取证小节（C7506/C7508 根因） | ch14-sm120-tma-ws.tex |
| Attention 数学、online softmax 流水、FA1/FA2/FA3 演进 | ch15-attn-math.tex |
| FA2（一）：Split-Q + MMA 多级流水、Casual mask | ch16-fa2-splitq-mma.tex |
| FA2（二）：TMA + WS 双流水、mbarrier 拓扑、坑七（setmaxnreg 约束） | ch17-fa2-tma-ws.tex |
| FA3：双 consumer warpgroup、寄存器再分配 | ch18-fa3-dual-consumer.tex |
| FFPA Split-D：大 head_dim（D>256）分块、两阶段 merge | ch19-ffpa-split-d.tex |
Part IV：CuTe（CUTLASS）（配图全部 TikZ 内联）
| 任务/症状 | 章 |
|---|---|
| CuTe 官方白皮书中文译注（第一手定义）：布局表示、张量、布局代数（拼接/合并/复合/补/切分/分块/逆） | ch19b-cute-whitepaper-zh.tex（分节在 chapters/wp/，wp6 另含 2 张官方架构截图） |
| Layout 基础与代数（colex/mode/compose/inverse/product/divide） | ch20-cute-layout.tex |
| Tensor 与 TiledCopy（thrval 引擎、g2s 分解） | ch21-cute-tensor-tiledcopy.tex |
| TiledMMA 与 fragment 布局 | ch22-cute-tiledmma.tex |
| Swizzle<B,M,S> 与 TMA copy | ch23-cute-swizzle-tma.tex |
| CuTe HGEMM 实战（kStage 流水、三级划分） | ch24-cute-hgemm.tex |
| CuTe FlashAttention 三实现对照 | ch25-cute-flash-attn.tex |
| CuTe FFPA Split-D 类型代数 | ch26-cute-ffpa.tex |
| SM120 持久化 FlashAttention（persist-D、WS 1+1、persistent CTA、scale 融合；全书唯一超越 cuDNN SDPA 的 attention kernel，240.1/230.2 = 1.04×） | ch26b-cute-persist-d-flash-attn.tex（书内第 28 章） |
| SM120 大 head_dim non-WS Split-D（tile 128×128、256T 全员 MMA + tid=0 内联 TMA、K/V stages 解耦：(3,2) D320=204.1T=2.93× cuDNN；ffpa-attn split_d 同源教学集，含寄存器 spill 机理） | ch26c-cute-split-d-sm120.tex（书内第 29 章） |
CuTe 进阶深度参考（colfax / cute-zhihu，独立文档）
当 ch20-26 的讲解不够用时，按问题深度依次下钻：
| 需求 | 位置 |
|---|---|
| 布局代数的第一手定义与动机（NVIDIA 官方白皮书，含 PyCuTe 参考实现） | 主书 ch19b（chapters/wp/wp0-7.tex） |
| 布局代数的数学严格化：范畴论视角（Tuple/Nest 范畴、可处理布局、复合/逻辑乘积/逻辑切分定理）+ 范畴论入门 | colfax/sec/cf-ch2*.tex、cf-ch3*.tex、cf-ch4*.tex、cf-appa.tex |
| reed 系列 13 篇：Layout → 代数与几何解释 → Tensor → Copy/MMA 抽象 → Swizzle → Hopper/mbarrier/TMA → 21bit TMA descriptor → simple/pipeline/efficient GEMM | cute-zhihu/chapters/reed-01~13-*.tex |
| 竹熙佳处系列 7 篇：tiled copy / tiled mma / compose & inverse / product & divide / TMA copy / permutationMNK / async pipeline | cute-zhihu/chapters/zhuxi-01~07-*.tex |
| 知乎文章全文快速 grep（7 篇已落 markdown） | references/fulltext/*.md |
查阅顺序建议：概念不清先读主书章 → 定义有疑义查 ch19b 白皮书 → 需要数学
严格性（定理证明/范畴论）查 colfax → 想看教学式推导与作者视角查 cute-zhihu。
Part V：FP8/FP4 HGEMM 篇（量化矩阵乘，PDF 打印章号 30-33；配图全部 TikZ 内联）
| 任务/症状 | 章 |
|---|---|
| e4m3 格式、per-row/per-block 量化、量化前处理 kernel（A/B 两侧量化链逐 bit） | ch34-fp8-quant-gemm.tex |
| FP8 CuTe 主 GEMM、在线反量化、cuBLAS BF16 对比、tile 档位扫描 | ch35-fp8-gemm-cute.tex |
| NVFP4 线格式（e2m1 数据 + ue4m3 SF）、512 B SF 原子 gmem 布局、两级折叠恒等式、误差模型 | ch36-fp4-gemm-quant.tex |
block-scaled MMA（mma.sync...m16n8k64.kind::mxf4nvf4.scale_vec::4X）、SF smem 流水、workspace 六段布局、形状自适应 tile、cuBLAS 对比 | ch37-fp4-gemm-cute.tex |
教学源码在 LeetCUDA 本仓 fp8_gemm.cuh / fp4_gemm.cuh（在线量化 GEMM，非
ffpa-attn 仓）；bench 入口 ./bin/leetcuda_bench_sm120a.bin --fp8-gemm <N> /
--fp4-gemm <N> / --fp4-gemm-sweep <N>。
Part VI：FP8/FP4 Attention 篇（量化注意力，PDF 打印章号 34-40，文件名 ch27-33；配图全部 TikZ 内联）
| 任务/症状 | 章 |
|---|---|
| 浮点位域、量化格点、ESS 分解等数学基础 | ch27-quant-attn-math.tex |
| FP8 量化前处理链（per-tensor/per-block/per-channel） | ch28-fp8-quant-aux.tex |
| FP8 persist-D 主 kernel、scale 折叠、reorg-free | ch29-fp8-persist-d.tex |
| FP8 split-D 与 M4N2 TiledMMA（大 D 两堵墙） | ch30-fp8-split-d-m4n2.tex |
| NVFP4 格式、量化链、KV perm32 置换 | ch31-fp4-nvfp4-quant.tex |
| FP4 persist-D、两级 P 量化 | ch32-fp4-persist-d.tex |
| FP8/FP4 bench 矩阵、精度方法论、竞品对照 | ch33-fp8-fp4-bench.tex |
配套生产级参考：ffpa-attn 仓库。 Part VI 各章的教学 kernel 均源自 ffpa-attn（书锚定其某个 commit 的冻结基线）， 而 ffpa-attn 本身是 sm_120（Blackwell）下 CuTe fp16/fp8/fp4 attention 算子的 峰值性能实现，可作为高质量 attention 算子参考库直接使用：
- fp16：超过 FA-2 标准实现；
- fp8：持平甚至略优于 SageAttention2；
- fp4（NVFP4）：明显优于 SageAttention3；
- large head_dim 专门优化：split-D / M4N2 TiledMMA 家族覆盖 D>256 大 D 场景 （教学对照 ch19/ch26/ch30）。
其余生产级沉淀：native/CuTe kernel 家族与特性矩阵、NHD 布局零拷贝（packed 与 strided 通用）、量化数学与 scale 折叠、性能 RFC 与已证伪清单、bench 与精度验证 方法论。深度任务应书章 + ffpa-attn 源码/文档联合使用：
- 仓库定位（记为
<FFPA_ATTN_DIR>）：环境变量FFPA_ATTN_DIR→ 当前目录及 父/兄弟目录查找ffpa-attn/→~/ffpa-attn→ 均无则git clone --depth 1 https://github.com/xlite-dev/ffpa-attn.git - 优先联合
<FFPA_ATTN_DIR>/.github/skills/ffpa-cuda-understandskill： 该 skill 是 ffpa-attn CUDA 后端的全景知识库（架构分发链路、kernel 家族、 量化数学原理、RFC 进度与验证方法论），与本书的 Part V/VI 章节互补——书讲 原理推导与教学实现，ffpa-cuda-understand 讲生产实现与工程决策。
附录（快速事实查询）
| 需求 | 附录 |
|---|---|
| common.cuh 逐段解析（TMA/mbarrier/setmaxnreg/WGMMA 宏封装、swizzle） | appendices/appA-common-toolbox.tex |
全书性能数据汇总（各卡 TFLOPS 基线表；ch11/12/14/16/17/18/24/25 章内另有 tab:chNN-perf 实测表） | appendices/appB-perf-data.tex |
| 构建指南：build.sh 架构×宏矩阵、-arch 目标选择、book/tests 用法、setmaxnreg 保留条件 | appendices/appC-build-guide.tex |
| 源码索引：kernel 源文件+行号定位表 | appendices/appD-source-index.tex |
| 参考文献与延伸阅读（含 PTX ISA 章节映射） | appendices/appE-references.tex |
Show full SKILL.md (434 more words)Show less
代码文件 references（kernels/interview 源码直定位）
按任务路由到章后，配套可参考源码直接按此表打开（行数为实测值；行号级
冻结映射见 appendices/appD-source-index.tex，老文件基准 commit 6cd32de；
fp8_gemm.cuh/fp4_gemm.cuh 及 bench/ 下 10 个 .cu（bench_leetcuda.cu +
9 模块 .cu，2026-09-29 迁入）的行号冻结以 book/scripts/anchors.yaml 为准）：
文件（相对 <LeetCUDA_DIR>/kernels/interview/） | 行数 | 内容 | 对应章 |
|---|---|---|---|
base.cuh | 909 | 架构/Roofline 速查、warp/block 归约与 dot、向量化与原子操作、softmax 三级递进、merge_attn_states、RMS/LayerNorm、RoPE 与转置 | ch01–07 |
sgemv.cuh | 102 | SGEMV 三种划分（warp-per-row K32/K128/K16） | ch08 |
sgemm.cuh | 434 | SGEMM 阶梯（block-tile/Vec4/双缓冲 → TF32 WMMA） | ch09–10 |
hgemm.cuh | 2100 | mma.sync m16n8k16 与 ldmatrix、XOR swizzle、TMA/WGMMA/mbarrier、SM120 TMA+WS、CuTe 对照片段与 CuTe HGEMM 实战 | ch11–14、ch20–21、ch23–24 |
flash_attn.cuh | 4179 | FA 原理头注释、FA2 split-Q+MMA、FA2 TMA+WS、FA3 双 consumer、CuTe FA 三实现、persist-D（flash_attn_cute_persist_d_sm120 L3635 起） | ch15–18、ch25、ch26b |
ffpa_attn.cuh | 641 | FFPA Split-D 全篇、双 TiledMMA/traits（M4N2 对照） | ch19、ch22、ch26 |
fp8_gemm.cuh | 958 | e4m3 量化前处理（per-row/per-block，L45-284）+ CuTe 主 GEMM、在线反量化（L303-955） | ch34-35 |
fp4_gemm.cuh | 1300 | NVFP4 在线量化 GEMM：A/B 两级量化 kernel 与误差模型（L58-366）+ block-scaled MMA 主 kernel、SF 流水、workspace（L395-1299）；形状自适应入口 fp4_gemm_use_wide_tile（MNK≥4096 用 128×256/s6，否则保守 128×128/s4） | ch36-37 |
common.cuh | 803 | TMA/mbarrier/setmaxnreg/WGMMA 宏封装与 swizzle 工具箱（逐段解析见 appA） | ch12–14、appA |
bench/bench_leetcuda.cu | 440 | main+CLI 总装（多 TU）：extern 声明 + 模板实例化 wrapper 调用；host test/bench 拆入 bench/{base,sgemv,sgemm,hgemm,fp8_gemm,fp4_gemm,flash_attn,ffpa_attn}.cu + bench/utils.cu 共享符号（~30 kernel 的 WHY+HOW 注释与 10 Phase 递进在 .cuh 内）；--bench 辅助工具与 --fp8-gemm/--fp4-gemm/--fp4-gemm-sweep GEMM 专项 bench CLI 不变（bin/leetcuda_bench_*.bin 源，build.sh --arch sm_120a --jobs 8 并行编译） | 全书总装、appC |
bench/deprecated/bench_attn.cu | — | （已弃用）FA2 CuTe TMA+MMA+WS vs cuDNN SDPA 专项 bench | ch00、ch17/25 |
bench/deprecated/bench_ffpa.cu | — | （已弃用）FFPA Split-D attention 专项 bench | ch19、ch33 |
bench/deprecated/bench_sgemm.cu | — | （已弃用）sgemm.cuh 全 kernel 性能+精度 bench | ch09–12 |
bench/deprecated/bench_sdpa.py | — | （已弃用）PyTorch SDPA 参照计时 | ch00、appB |
book/tests/chNN_*.cu | — | 每章最小正确性测试（CPU fp64 对拍），build_tests.sh --arch sm_120a --all | 各章 |
Part VI（ch27–33，量化 attention）教学 kernel 源码在 <FFPA_ATTN_DIR>/csrc/cuffpa/
（appD 内附逐章 GitHub permalink）；Part V（ch34–37，量化 GEMM）教学 kernel 源码在
本仓 kernels/interview/fp8_gemm.cuh 与 fp4_gemm.cuh。
第三步：按需读取纪律
- 先路由后读取：一次只读任务相关的 1-2 章；章内先看
\section/\subsection目录再定位行区间（每章 500-900 行，全文读入浪费上下文）。 - 源码交叉：书按行号引用
<LeetCUDA_DIR>/kernels/interview/*.cuh冻结基线 （common/sgemv/sgemm/hgemm/flash_attn/ffpa_attn/fp8_gemm/fp4_gemm.cuh），需要完整实现时用appendices/appD-source-index.tex定位再读源文件。 - 图：章内图全部为 TikZ inline（
tikzpicture就在章节 tex 内）。看图：pdftotext book.pdf -按图题关键词定位页码 →pdftoppm -r 150 -png整页 渲染 → view 工具目检。改图：直接改 tex 内tikzpicture，改后必须真实渲染 （standalone 抽图或重建 PDF）+ view 视觉验收无文文/图文重叠（铁律与配方见 write-leetcuda-book skill 图管线）。figures/drawio/、figures/ffpa/下 png 为历史留档，仅作旧版对照。 - 验证 kernel 正确性：优先跑
book/tests/（CPU fp64 对拍，无 cuBLAS 依赖）； 性能验收用bash build.sh --arch sm_120a && ./bin/leetcuda_bench_sm120a.bin --bench --bhnd 1,32,8192,128 [--bench-all]；FP8/FP4 GEMM 专项用--fp8-gemm <N>/--fp4-gemm <N>（形状自适应行）/--fp4-gemm-sweep <N>（tile 档位扫描） （产物落bin/，bench 辅助工具源在kernels/interview/bench/；构建细节见 appC）。 - PDF 兜底：tex 源不适合读时（如只要结论），
pdftotext book.pdf -按页抽取； colfax 与 cute-zhihu 两份独立文档同理（pdftotext colfax/colfax-cute-zh.pdf -、pdftotext cute-zhihu/cute-zhihu.pdf -）。 - 不要修改书的内容除非任务明确要求；本 skill 与书保持同步演进。
已知重要结论速查（细节见对应章节）
- setmaxnreg 被 ptxas 静默丢弃的真根因：
cp.async.bulk.tensor目的地址写shared::cluster会被视为 implicit extern 边界（C7506）；写shared::cta则 sm_120a/sm_120f 均保留。另需__launch_bounds__(N, 1)（否则 C7508）。取证过程 见 ch14 专门小节；SASS 助记符是USETMAXREG（grepsetmaxnreg查 SASS 会假阴性）。 - bench A/B 口径：同轮 SDPA 参照一致性、min-of-N、PSNR/Max-Err 双指标——见 ch00 与 appB。
- GPU 规格口径（2026-09-22 运行时实测核准）：PRO 5000 = 110 SM
（
torch.cuda.get_device_properties().multi_processor_count）；历史误传 "96 SM" 已全书清除（96 系与 Full GB202 裸 die 的 TPC 数混淆）；5090 = 170 SM、 PRO 6000 = 188 SM。硬件规格一律以运行时实测为准，书内不写其他卡的 SM 数。 - 性能基线速查（PRO 5000，2026-09-22
--bench --bench-all）：ch26b persist-D 240.1 TFLOPS = 1.04× cuDNN SDPA（230.2）；HGEMM CuTe Swizzle 追平/超越 cuBLAS （最高 1.08× f16、1.48× f32）；明细见各章tab:ch11/12/14/16/17/18/24/25-perf。 - FP8/FP4 GEMM 基线（PRO 5000，cuBLAS BF16 参照 163.1/163.4 T @ 4096³/8192³）： FP8 ws 128×256/s2 kernel-only 443.7/487.8 T（2.72×/2.99×），e2e 含量化 186.5/314.6 T；NVFP4 形状自适应（MNK≥4096 宽档 128×256/s6，否则保守 128×128/s4）：4096³ kernel-only 603.4/614.1 T（3.69×/3.76×）、e2e 408.2/414.5 T 反超 FP8 e2e；8192³ 639.6/648.6 T、e2e 467.4/472.3；B 离线部署口径 550.5/561.2 T（= kernel 的 91-93%）。量化成本 109 µs 是加性的，宽档下占 e2e 32%——占比上升是主 kernel 变快所致，不是量化变慢（ch37 有防误读句）。
© xlite-dev, GPL-3.0. Rendered from Markdown: HTML in the file is shown as text, images as links, and headings moved down two levels. Raw file
Files
Just SKILL.md in kernels/interview/book/skills/leetcuda-cpp-kernel of xlite-dev/LeetCUDA.
Open the folder on GitHubat commit c5e985f
Compare with similar skills
Leetcuda Cpp Kernel next to the 5 skills that share the most tags, products or categories with it. Stars are the repository's; “used in” counts other GitHub owners with a copy.
| Skill | Stars | Used in | Tokens | Auto-check | Licence | Repo updated |
|---|---|---|---|---|---|---|
| Leetcuda Cpp Kernel this skillxlite-dev/LeetCUDA | 12k | — | ~3.5k | Automated safety check: Pass | GPL-3.0 | |
| Cuda Cpp Kernelvipshop/cache-dit | 1.3k | — | ~2.3k | Automated safety check: Pass | Apache-2.0 | |
| GPU OptimizationspiriMirror/libuipc | 336 | — | ~3.6k | Automated safety check: Pass | Apache-2.0 | |
| Add Jit Kernelguqiong96/Lsglang | 144 | 1 repos | ~10k | Automated safety check: Pass | Apache-2.0 | |
| At Dispatch V2intel/torch-xpu-ops | 115 | 3 repos | ~2.2k | Automated safety check: Pass | Apache-2.0 | |
| Add Jit Kernelsgl-project/sglang | 37k | — | ~13k | Automated safety check: Pass | Apache-2.0 |
Similar skills
GPU Optimization
spiriMirror/libuipc
GPU optimization workflow using uipc.profile, uipc.profile.nsight, and Nsight Compute CLI.
Add Jit Kernel
guqiong96/Lsglang
Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang's jitkernel module
- Official
At Dispatch V2
intel/torch-xpu-ops
Convert PyTorch ATDISPATCH macros to ATDISPATCHV2 format in ATen C++ code.
Add Jit Kernel
sgl-project/sglang
Step-by-step tutorial for adding a new lightweight JIT CUDA kernel to sglang.kernels JIT infrastructure and public operator groups
Cuda
mohitmishra786/low-level-dev-skills
CUDA C/C++ skill for NVIDIA GPU kernel programming. An agent skill from mohitmishra786/low-level-dev-skills.
More from xlite-dev/LeetCUDA
Write Leetcuda Book
xlite-dev/LeetCUDA
把 LeetCUDA kernels/interview 教学源码（base/sgemv/sgemm/hgemm/flashattn/ffpaattn/fp8gemm.cuh/fp4gemm.cuh）+ ffpa-attn CuTe sm120 源码（csrc/cuffpa/cute fp8/fp4，commit 861d75e）写成中文 CUDA 技术书（6 Part 38 章 + CuTe…
Related
Categories
Questions about Leetcuda Cpp Kernel
What does Leetcuda Cpp Kernel do?
LeetCUDA 中文技术书（584 页，XeLaTeX 源）按需查阅 skill——写、优化、调试或 review CUDA C++/PTX kernel 时的权威参考路由层。当任务涉及：GPU 架构/Roofline/ occupancy、向量化与 coalescing、warp/block reduce、softmax（online/LSE merge）、…. Leetcuda Cpp Kernel is an agent skill from xlite-dev/LeetCUDA.
When should I use Leetcuda Cpp Kernel?
Leetcuda Cpp Kernel fits situations like: tasks that involve GPU and accelerator computing; tasks that involve CSS and styling; tasks that involve LaTeX.
How do I install Leetcuda Cpp Kernel in Claude Code?
Run `npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a claude-code`. Or copy the skill folder (kernels/interview/book/skills/leetcuda-cpp-kernel in xlite-dev/LeetCUDA) into .claude/skills/leetcuda-cpp-kernel in your project. Claude Code loads it when a task matches its description.
How do I install Leetcuda Cpp Kernel in Codex?
Run `npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a codex`. Or copy the skill folder (kernels/interview/book/skills/leetcuda-cpp-kernel in xlite-dev/LeetCUDA) into .agents/skills/leetcuda-cpp-kernel in your project. Codex loads it when a task matches its description.
Can I use Leetcuda Cpp Kernel in Cursor, Gemini CLI or GitHub Copilot?
Cursor, Gemini CLI, GitHub Copilot and OpenCode also load SKILL.md folders. With the skills CLI, run `npx skills add xlite-dev/LeetCUDA --skill leetcuda-cpp-kernel -a cursor` (or -a gemini-cli, github-copilot or opencode for the others). To copy it by hand, put the folder in .cursor/skills/leetcuda-cpp-kernel, .gemini/skills/leetcuda-cpp-kernel, .github/skills/leetcuda-cpp-kernel and .opencode/skills/leetcuda-cpp-kernel in your project.
What does Leetcuda Cpp Kernel need to run?
Going by SKILL.md and its folder, Leetcuda Cpp Kernel needs the command-line tools its instructions call (pdftotext, git and pdftoppm).
Does Leetcuda Cpp Kernel access the network?
SKILL.md names 1 domain. In commands or code: github.com; the agent is likely to contact it when it follows the instructions. This is read from the text; nothing was executed.
Is Leetcuda Cpp Kernel safe to install?
Our automated static check of SKILL.md found no risky patterns, such as piping downloads into a shell, reading credential files or hidden Unicode. It is not a guarantee. Review the folder before installing.
What licence does Leetcuda Cpp Kernel use?
Leetcuda Cpp Kernel is published under the GPL-3.0 licence (the repository's licence). It allows redistribution, so the full SKILL.md is shown on this page.
How many tokens does Leetcuda Cpp Kernel use?
About 3.5k tokens (SKILL.md is roughly 14k characters). Agents keep only the skill's name and description in context until a task matches; then they load SKILL.md in full.
What are the alternatives to Leetcuda Cpp Kernel?
Skills that share tags, products or a category with Leetcuda Cpp Kernel: Cuda Cpp Kernel (vipshop/cache-dit, 1.3k stars), GPU Optimization (spiriMirror/libuipc, 336 stars), Add Jit Kernel (guqiong96/Lsglang, 144 stars) and At Dispatch V2 (intel/torch-xpu-ops, 115 stars). The comparison table on this page puts their stars, adoption, token cost, safety result and licence side by side.
Who maintains Leetcuda Cpp Kernel?
xlite-dev (a GitHub organization) maintains it in xlite-dev/LeetCUDA, which has 12,035 GitHub stars. The repository holds 2 skills in this directory. The repository was last updated on September 30, 2026.
Source: xlite-dev/LeetCUDA on GitHub. Facts on this page come from the repository at the commit we read; the author's words are quoted as theirs.