GPU MODE

The single most important community for this roadmap.

What it is

GPU MODE (formerly CUDA MODE) is a Discord community + YouTube lecture series + kernel-writing leaderboard, founded by Mark Saroufim (Meta) and Andreas Koepf. It is where the people who write FlashAttention, Marlin, ThunderKittens, vLLM kernels, Liger, Unsloth, TorchTitan actually hang out. Guest lectures are given by the author of the thing you just read the paper on.

If you internalize one thing from this file: join the Discord in Month 1. Don’t wait until you “feel ready.”

Discord channel map (as of 2026 organization)

Don’t spam the general channels. Route your posts:

  • #introduce-yourself — one message, once. Say who you are, what you’re building, what GPU you have.

  • #general — low-signal chatter. Skim, don’t rely on.

  • #triton, #cutlass, #cuda-beginners, #pmpp-book — topic-specific. Ask questions here after you’ve RTFM’d.

  • #lectures — discussion around each YouTube lecture. Post specific questions with timestamps.

  • #kernel-challenges / leaderboard channels — competitions. This is where you build reputation.

  • #help — for genuine blockers. Include: what you’re trying, what happens, minimal repro, nvidia-smi output.

  • #jobs — lurk for hiring signals. Do not post “looking for a job.” Let your kernel PRs be your resume.

The lecture recommendations (~40 min each, ordered for THIS roadmap)

All live at https://github.com/gpu-mode/lectures with code + slides. Watch these on rotation — one per week is a good pace. Priority-ordered for an inference engineer trajectory:

Phase 2 (GPU arch + kernels)

  1. Lec 1 — Profiling & Integrating CUDA in PyTorch (Mark Saroufim). Start here. The Python↔CUDA bridge you’ll use in every project.

  2. Lec 4 — Compute and Memory Architecture (Thomas Viehmann). The hardware model, right.

  3. Lec 3 + Lec 5 — Getting Started With CUDA + Going Further (Jeremy Howard). Zero-to-something in the tool.

  4. Lec 8 — CUDA Performance Checklist (Mark Saroufim). The checklist you should tape to your monitor.

  5. Lec 9 — Reductions (Mark Saroufim). Where you first meet warp shuffles.

  6. Lec 14 — Practitioner’s Guide to Triton (Umer Adil). The pragmatic entry point.

  7. Lec 29 — Triton Internals (Kapil Sharma). How the compiler works under you.

  8. Lec 16 — Hands-on profiling (Taylor Robbie). Nsight in practice.

Phase 3 (attention + fusion)

  1. Lec 12 — Flash Attention (Thomas Viehmann). Companion to reading FA1/FA2.

  2. Lec 36 — CUTLASS and Flash Attention 3 (Jay Shah). Hopper-era FA internals.

  3. Lec 18 — Fused Kernels (Kapil Sharma).

  4. Lec 23 — Tensor Cores (Vijay Thakkar & Pradeep Ramani — the CUTLASS core devs).

  5. Lec 15 — CUTLASS (Eric Auld). The template-mountain intro.

Phase 4 (engines)

  1. Lec 22 — Speculative Decoding in vLLM (Cade Daniel). Straight from the vLLM contributor.

  2. Lec 35 — SGLang Performance Optimization (Yineng Zhang). Straight from SGLang core dev.

  3. Lec 40 — FlashInfer (Zihao Ye). The author of the library.

  4. Lec 39 — TorchTitan (Mark Saroufim, Tianyu Liu). Modern training stack.

Phase 5 (quantization)

  1. Lec 7 — Advanced Quantization (Charles Hernandez).

  2. Lec 30 — Quantized training (Thien Tran).

  3. Lec 33 — BitBLAS (Wang Lei). W4A16/W2A16 GEMM.

  4. Lec 34 — Low Bit Triton Kernels (Hicham Badri, HQQ author).

  5. Lec 32 — Unsloth (Daniel Han). LLM systems engineering war stories.

Phase 6 (distributed)

  1. Lec 17 — GPU Collective Communication / NCCL (Dan Johnson).

  2. Lec 67 — NCCL & NVSHMEM (Jeff Hammond).

  3. Lec 13 — Ring Attention (Andreas Koepf).

  4. Lec 78 — Iris: Multi-GPU Triton.

Awareness (later, or as needed)

  • Lec 25 — Composable Kernel (AMD/CK stack).

  • Lec 26 — SYCL MODE (Intel).

  • Lec 31 — Metal Kernels (Nikita Shulga).

  • Lec 37 — SASS & GPU Microarchitecture (Arun Demeure). Deep hardware.

  • Lec 84 — Numerics and AI (Paulius Micikevicius — the FP8 spec author).

  • Lec 103 — CuTe Layout Algebra (Jack Carlisle & Jay Shah). Category theory / advanced.

  • Lec 104 — Gluon: Tile-Based GPU Programming (Triton core devs).

  • Lec 41 — CUDA Docs for Humans (Charles Frye). Reference for reference.

Leaderboard participation strategy

GPU MODE runs periodic kernel-writing competitions (KernelBench-style problems: “implement this op faster than the reference on this hardware”). This is the highest signal-to-noise reputation-building mechanism in the community.

Strategy:

  1. Skip the first competition you see. Just read the writeups from the winners. Study what techniques won.

  2. Enter competition #2 with a modest goal. Aim for top 50%. Ship a working solution. Learn the harness.

  3. Enter competition #3 to place. By now you know the profiler workflow, the numerical tolerance harness, the common tricks. Aim top 10%.

  4. Winning is not the goal. Learning is. Even a mid-place solution + a good writeup is more career signal than 90% of Kaggle finishes.

Watch the winners’ writeup posts — they routinely become the seed of a paper or a merged PR to vLLM/CUTLASS. That’s the pipeline you’re trying to enter.

“How to speak GPU-MODE” primer

The dialect. Learn these terms before your first post so you don’t sound like a tourist:

  • “achieved bandwidth” — measured GB/s of your kernel. Always report vs. peak.

  • “roofline” — the compute-vs-bandwidth diagram. Every perf discussion assumes this.

  • “arith intensity” / “AI” — FLOPs/byte. Below the ridge = memory-bound.

  • “NCU” / “Nsight Compute” — the per-kernel profiler. Report NCU screenshots with claims.

  • “Nsys” / “Nsight Systems” — the timeline profiler.

  • “warp / block / cluster” — 32 threads / thread block / (Hopper+) thread-block cluster.

  • “SMEM” / “shared memory” / “scratchpad” — the programmer-managed cache.

  • “cp.async” / “TMA” — async HBM→SMEM copy (Ampere / Hopper). Different generations, don’t conflate.

  • “WGMMA” / “tcgen05” — async warp-group tensor-core MMA (Hopper / Blackwell).

  • “SOL” — “speed of light,” the theoretical peak for the op.

  • “pingpong” — double-buffered producer/consumer warp specialization.

  • “MMA / HMMA / QMMA” — tensor-core matrix-multiply-accumulate instructions.

  • “nvcc” / “ptx” / “sass” — source → PTX (virtual ISA) → SASS (hardware ISA). Real perf debugging goes to SASS.

  • “eager” vs “compiled” vs “graph-captured” — the three PyTorch execution modes. Know which you’re benchmarking.

If a term appears in Discord that you don’t recognize, screenshot → search → close the tab. Do this 20 times and you’re fluent.

First-message template (post in #introduce-yourself, once)

Hey — [Name] here. Applied ML eng by day (recsys / agentic services at $CO), spending nights learning kernels + inference internals with the goal of contributing to vLLM/SGLang and eventually working on inference infra. Currently on Phase 2 of a self-designed roadmap: SGEMM ladder on my 3090 [or: renting H100 via RunPod for FA3 work]. Plan to lurk mostly; will surface with benchmark writeups. Happy to help anyone on the PMPP exercises — that’s where I am.

One message. Then quiet. Contribute value, don’t announce intentions.

First contribution ideas (Phase 2–4)

  • Reproduce a leaderboard-winning kernel and post your NCU comparison.

  • Take one lecture you watched and post a corrected/annotated notebook.

  • Cross-post your Simon-Boehm-inspired SGEMM ladder with your card’s numbers.

  • Publish your quant bake-off (Phase 5) with a table — that’s exactly the content this community boosts.

  • File a good issue in the vLLM/SGLang repo (see 03_open_source_engines.md) and mention it in #general — not for karma, but because someone here likely has the answer.

What NOT to do

  • Don’t @ Tri Dao, Horace He, Mark Saroufim, or any named contributor unsolicited. Ever.

  • Don’t ask “can someone teach me CUDA?” Read PMPP. Do the exercises. Then ask specific questions.

  • Don’t share code without a benchmark. “I think this is faster” is not a claim.

  • Don’t LLM-generate a wall of code and ask why it doesn’t work. You will be ignored.

  • Don’t post crypto-style hype. This community has an extremely low tolerance for it.