Rung 2 — CUDA SGEMM Ladder¶

Aligned phase: Phase 2 (GPU arch + CUDA + Triton) Ship by: end of M4 Effort: ~4 weeks, ~40 hrs total Signal: mid. This is the first serious rung — it proves you can write CUDA.

The artifact spec¶

Write a series of CUDA SGEMM kernels for FP32 C = A @ B on 4096×4096 matrices, iterating from naive to within striking distance of cuBLAS. Each kernel is a separate .cu file and a separate row in the benchmark table.

Version	Optimization	Target % of cuBLAS
v0	Naive one-thread-per-C-element	~1-2%
v1	Global memory coalescing	~5-10%
v2	Shared memory blocking	~15-25%
v3	1D thread tile (each thread computes 8 elements)	~30-40%
v4	2D thread tile (8×8 per thread)	~50-60%
v5	Vectorized loads (`float4`)	~65-75%
v6	Warp tiling + double buffering	~80-90%
v7 (stretch)	WMMA on tensor cores (drop to FP16 or TF32)	~95%+ (different game)

Reference: Simon Boehm’s CUDA-MMM series is your literal target. Read all 10 blog posts before starting. Then reimplement from memory. That’s the exercise.

Benchmark harness:

Fixed clocks: nvidia-smi -lgc <base_clock> and nvidia-smi -lmc <mem_clock>.
10× warmup, 100× measured, report median GFLOPS and P99.
Verify correctness against cuBLAS output (max abs diff < 1e-3).
Run on your 3090 (or rent a 4090 for one afternoon to make the numbers look better on the plot).

Repo structure¶

sgemm-ladder/
├── README.md              # Blog-quality writeup with hero plot
├── kernels/
│   ├── v0_naive.cu
│   ├── v1_coalesce.cu
│   ├── v2_smem.cu
│   ├── v3_1dtile.cu
│   ├── v4_2dtile.cu
│   ├── v5_vectorize.cu
│   ├── v6_warptile.cu
│   └── v7_wmma.cu         # optional
├── bench/
│   ├── run.py             # sweeps kernels, writes CSV
│   ├── plot.py            # regenerates hero plot
│   └── pin_clocks.sh
├── results/
│   ├── raw.csv
│   ├── hero.png           # GFLOPS bar chart + % cuBLAS
│   └── roofline.png       # each kernel on the 3090 roofline
├── src/
│   ├── bench.cu           # unified harness
│   └── reference.cu       # cuBLAS baseline
└── Makefile

Where to post¶

GitHub: github.com/<you>/sgemm-ladder. Pin the repo.
Blog: cross-post README. Title: “CUDA SGEMM: from 0.5% to 89% of cuBLAS on a 3090” — the % number in the title is the hook.
r/MachineLearning + r/CUDA: submit link. Sunday morning ET.
GPU MODE Discord #cuda: share and ask for critique.
Twitter/X: thread the top 3 optimization insights with the hero plot. Tag @marksaroufim @cHHillee (do not @-spam; only if the content actually stands).
HN Show HN: worth trying now.
LinkedIn: skip.

What signals it sends¶

“I can write CUDA that is not embarrassingly slow.”
“I can reason about coalescing, occupancy, shared memory, warp tiling.”
“I have used ncu/nsys and can read a profiling report.”
“I can iterate 8 versions of a kernel without giving up.”

This is the artifact that unlocks the GPU MODE Discord conversation. With rung 2 in hand you can DM real kernel engineers a specific question and they will engage. Without it, they will politely ignore you.

Past examples¶

Simon Boehm CUDA-MMM — https://siboehm.com/articles/22/CUDA-MMM. Your literal reference. Cite it explicitly in your README.
Colfax Research SGEMM writeup — Google “Colfax CUDA SGEMM.” The professional version.
NVIDIA CUTLASS docs — https://github.com/NVIDIA/cutlass. Overkill but read the CuTe quickstart section for context.
Lei Mao’s CUDA blog — https://leimao.github.io/tags/CUDA/ has multiple kernel walkthroughs at your target level.

Common mistakes¶

Cargo-culting Simon Boehm. If you copy code without understanding, you learn nothing. Do it from memory after reading.
Not profiling with ncu. Every optimization should be justified with a metric change: DRAM throughput up, L1 hits up, occupancy at X%, etc. Screenshots of ncu in the blog post are gold.
Ignoring TF32. On Ampere+, cuBLAS uses TF32 for FP32 matmul by default. Compare apples to apples: either force cuBLAS to strict FP32 (cublasSetMathMode(handle, CUBLAS_PEDANTIC_MATH)) or state that both use TF32.
Comparing on the wrong shape. 4096×4096 is the standard. Do NOT benchmark on 1024×1024 where launch overhead dominates.
Not showing failures. If v6 is slower than v5 on your card, include it and explain why. Honest > polished.

Success criteria¶

Final version ≥60% of cuBLAS FP32 on your card
Hero plot in repo + blog
ncu metrics table for at least 3 versions
Repo has ≥25 stars within 2 weeks of posting (soft target)
At least one person on GPU MODE engages substantively
You can whiteboard a coalesced kernel from scratch in an study

If ≥5 of 6 check, ship rung 3.