Kernel Dev 训练营/ 提示与讲解
提示与讲解

Kernel Dev 训练营 · 提示与讲解

每道题一页:背景、参考解逐段讲解、图解、类比和钉死版本的源码链接。建议先自己写,卡住了再来看。

Track C · CUDA C++

硬件映射:线程索引、合并访存、shared memory、warp 原语、数据复用。

C1Vector add + 融合 bias-GELU:索引、launch 配置、fusion35 min · 易
C2Reduction:warp shuffle + shared memory + atomic45 min · 中
C3矩阵转置:coalescing 与 bank conflict40 min · 中
C4GEMM 三连:naive → shared-memory tiling → register tiling90 min · 难

Track T · Triton

block 级编程、mask、融合、reduction、tl.dot、autotune、FlashAttention。

T1Triton 入门:融合 add + GELU25 min · 易
T2融合 row softmax:数值稳定 + 一次读写35 min · 中
T3LayerNorm / RMSNorm forward40 min · 中
T4Triton GEMM:tl.dot、grouped ordering、autotune60 min · 中
T5FlashAttention forward:tiling + online softmax90 min · 难

Track L · 推理 attention 从零写

Paged KV cache、decode attention、GQA、split-KV:真实 serving 的形状。

L1Paged KV cache:slot mapping + 写 cache35 min · 易
L2Paged decode attention + GQA60 min · 难
L3Split-KV decode(Flash-Decoding)+ LSE 合并60 min · 难

Track F · FlashInfer

对着 FlashInfer 的真实 API 实现并逐数值对拍,最后进入它的源码和贡献流程。

F1FlashInfer 分页格式(CSR)+ plan / run decode40 min · 易
F2merge_state / merge_states:以 2 为底的 LSE 合并40 min · 中
F3fused_add_rmsnorm(原地)+ silu_and_mul35 min · 中
F4不排序的 top-k / top-p:对浮点位模式二分60 min · 难
F5FlashInfer 贡献实战:源码地图、JIT、测试、benchmark、PR 规范180 min · 中

Track G · Layout / CuTe

数据在 warp / lane / register 上怎么分布;CuTe 的布局代数。

G1Layout 计算器:BlockedLayout / SliceLayout 到线程的映射45 min · 中
G2CuTe 布局代数:coalesce、composition、complement、divide60 min · 难

Track M · MoE 与低精度

路由与对齐、grouped GEMM、FP8 block scaling:DeepSeek-V3 / DeepGEMM 的配方。

M1MoE 路由 + moe_align_block_size40 min · 中
M2Fused MoE grouped GEMM(vLLM fused_moe 结构)75 min · 难
M3FP8 细粒度量化 + block-scaled GEMM(DeepSeek-V3 / DeepGEMM 配方)60 min · 难

Track X · MLX / Metal

Apple silicon 上的自定义 kernel:elementwise、simdgroup 归约、simdgroup_matrix GEMM。

X1MLX 自定义 Metal kernel 入门:融合 add + GELU30 min · 易
X2Metal 归约:simdgroup + threadgroup 内存的 softmax / RMSNorm45 min · 中
X3Metal GEMM:threadgroup tiling → simdgroup_matrix75 min · 难

专题

跨题目的系统讲解。

专题CuTe 专题:用“布局代数”管住 kernel 里所有的下标
专题FlashInfer 原理:为 LLM serving 设计的 attention 引擎
专题Metal / MLX 与 CUDA / Triton:从硬件到 kernel 的对比
专题Tensor Core 与新指令集:从 WMMA 到 wgmma、TMA、tcgen05