monolith

Part 1: For you

The idea in one line

One persistent CUDA kernel runs every layer of a real language model, Qwen3-0.6B, on an RTX 5090. You make it about twice as fast on a gaming card, fold the last two helper kernels into it so that one launch produces each token, and build an X-ray that shows what each of the card's 170 SMs does during that token.

Why it's exciting

• One kernel per token. A megakernel is a single persistent CUDA kernel that runs the decode step without returning to the CPU. Upstream runs all 28 layers in one launch and then two small kernels for the LM head argmax; you fold those in, so each token is one launch with no idle gaps between kernels.
• The twist is in the barriers. AlpinDale's open-source kernel already decodes about 1,033 tokens per second at batch size 1. By its author's own measurement, though, "one-third of the per-layer time is synchronization overhead." Quantization shrinks the weight reads and leaves the barriers alone, so quantization by itself can't double the speed. The X-ray makes that visible.
• It's low-level for real. You dequantize FP8 and NVFP4 weights in registers with PTX cvt instructions, replace grid-wide barriers with atomic counters and release/acquire memory ordering, and stamp time with %globaltimer, clock64, and %smid.
• Nobody has shipped this. The research found no quantized (FP8 or NVFP4) megakernel for Qwen3-0.6B on the RTX 5090. Cohere writes, "We plan to release an RTX megakernel soon," and an open Mirage feature request asks for a per-SM profiler whose output "feeds a viewer drawing the SM-by-time chart used in the decks": today that chart lives in slides. The window is short, so ship soon.
• You build your own instrument. Nsight Compute hardware counters are likely blocked on RunPod's unprivileged pods. In-kernel timestamps work anywhere, and they show something a counter summary can't: every block, every phase, one token.
• The numbers are honest. llama.cpp, vLLM, and Hugging Face run on the same GPU, side by side, and no speed number counts until that exact build passes a correctness gate.
• The whole budget fits on one line. A bf16 decode step reads about 1.19 GB of weights. FP8 cuts that to about 0.60 GB and NVFP4 to about 0.40 GB, so you can predict each version's speed by hand and then check the prediction against the X-ray.

What the demo looks like

• The race. Four token streams, from monolith, llama.cpp, vLLM, and Hugging Face, replay side by side at the timings recorded on the same RTX 5090.
• The X-ray. An SM × time Gantt chart of one token: one row per thread block, one color per phase, and hatched bars where a block waits at a barrier. You can zoom into a single layer and hold the pointer over a bar to see its phase, SM, and duration.
• A version switcher that steps through the kernels: bf16 upstream, then FP8, then FP8 with fewer barriers, then NVFP4. The barrier share visibly grows after FP8 and shrinks after the barrier work.
• A "where the time goes" budget bar for each version, which splits one token into weight reads, barrier waits, attention, and the LM head, next to the roofline minimum.
• A results table that puts quality (perplexity, greedy agreement, and KL divergence) next to speed, so a fast but broken variant can't hide.
• A short code tour of the two ideas that matter: the dequant inner loop and the counter wait that replaces a grid barrier.

How it works

1. Probe the pod. On a RunPod RTX 5090, check whether Nsight Compute counters work, how finely %globaltimer ticks, whether the FP4 cvt instruction exists on sm_120a, and how much DRAM bandwidth a plain streaming kernel achieves.
2. Reproduce the upstream kernel. Build AlpinDale's qwen_megakernel, match Hugging Face greedy tokens, and reproduce about 1,033 tokens per second. Measure llama.cpp, vLLM, and Hugging Face on the same pod.
3. Build X-ray. Every thread block writes a GPU timestamp at the start and end of each phase. An analyzer turns the dump into per-phase time, barrier wait, and idle time per SM.
4. Quantize to FP8. Pack the weights as FP8 with one scale per output row, and dequantize them in registers inside the matrix-vector products. Check quality against bf16.
5. Cut the barriers. Fuse phases, replace grid barriers with per-tile counters, retune the block count, prefetch during more of the waits, and fold the LM head into the kernel. Keep only what stays correct.
6. Try NVFP4. Pack the layers as 4-bit NVFP4 with block scales, keep sensitive matrices in FP8, and report quality honestly.
7. Measure everything once. Run every variant and baseline on one pod session, and export one recorded token trace per variant.
8. Ship it as a static web page on vm.ifkash.dev, with the kernel repo and a blog post.

Weekend plan

Cost

• GPU: 1× RTX 5090 on RunPod, $0.69 per hour on community cloud or $0.99 per hour on secure cloud. The fallback is an RTX PRO 6000 at $1.69 per hour. It has the same ISA but 188 SMs, so if you use it, you re-measure every number on that card.
• Time on the pod: about 15-25 hours, by estimate. CUDA work can't compile or run on the Mac, so the pod is up for most of the kernel work. Stop it while you write page code on the Mac.
• Total: about $15-25. The plan sets a hard stop at $40.
• Hosting: free. It's a static page on your VM, and it keeps working after the pod is gone.

What you have at the end

• A live link where anyone can race the token streams and scrub through an X-ray of every kernel version.
• The monolith kernel repo: FP8 and NVFP4 megakernels for Qwen3-0.6B, plus X-ray.
• A benchmark table of monolith against llama.cpp, vLLM, Hugging Face, and the upstream kernel, all on one RTX 5090, with quality next to speed.
• Recorded traces of one token for every version, which anyone can reload and inspect.
• A blog post titled something like "One kernel, one token, 170 SMs: making a megakernel twice as fast on a gaming GPU."

What might go wrong

Other ideas the research turned up

• DeepSelect. DeepSeek's TopK kernel behind its sparse attention (v1.0.0, MIT) runs 2-20× faster than torch.topk for top-k up to 4096. It builds only for sm_100a and sm_103a, so an sm_120 port with an animated explainer is open.
• Direct-P FP4 FlashAttention. Up to 2.13× over BF16 forward on GB200, but sm_100 only. A port to sm_120 block-scaled mma.sync is open.
• The missing Gated DeltaNet decode kernel. vLLM's fused kernel needs num_v_heads == 8*num_k_heads, so Qwen3.8-27B with multi-token prediction falls back to Triton on an RTX 5090, at 19.9 tok/s against 72.3 without it.
• GPU MODE's linear-algebra kernels. Batched QR, eigh, and Cholesky on hosted B200s. Whether the leaderboards still accept submissions is unverified.
• Faster Than Flash. Sparse flash decoding (ICML 2026), up to 11.6× at kernel level and 2.37× end to end over FA2 on an RTX 4090. The code has one commit.
• Tensor-core numerics on sm_120. Nobody has published how the RTX 5090's FP8 and NVFP4 mma.sync instructions accumulate and round. It's a cheap probe project.

Reading, if you want it

• AlpinDale's 5090 decode optimization post and its repo: the kernel this project forks, and the barrier numbers that motivate it.
• Look Ma, No Bubbles (Hazy Research): the on-GPU interpreter and counter-based sync that inspire the barrier work.
• Mirage Persistent Kernel: a compiler that generates megakernels, benchmarked on A100, H100, and B200.
• Cohere's megakernels post: a serving megakernel for a 30B-A3B MoE, and the RTX plans that make this project time-sensitive.
• AutoMegaKernel: agent-proposed megakernel schedules, including sm_120 numbers for bf16.
• KernelBench-Verified: why every speed number here passes a correctness gate first.
• Colfax's NVFP4 GEMM on RTX PRO 6000: block-scaled FP4 on sm_120 in detail.
• PTX ISA: the reference for cvt, release/acquire loads and stores, and the special registers X-ray reads.
• NVIDIA RTX Blackwell whitepaper: the RTX 5090's SMs, memory, and tensor cores.

---

Part 2: For the coding agent

Mission

Build monolith, a faster decode megakernel for Qwen3-0.6B at batch size 1 on the RTX 5090, and an instrument that shows where its time goes:

• Kernels: FP8 and NVFP4 weight-only megakernels, forked from AlpinDale's MIT-licensed qwen_megakernel, with dequantization in registers and fewer grid-wide barriers.
• X-ray: A per-block phase timestamp tracer behind a compile-time flag, plus an analyzer and a trace exporter.
• Harness: A benchmark and quality harness that measures every variant and every baseline (upstream, llama.cpp, vLLM, and Hugging Face) on the same pod.
• Demo page: A static page that replays recorded traces and token streams: the race, the X-ray Gantt chart, the budget bar, the results table, and a code tour.

The final artifact is a static website with no inference server; the page needs no GPU. The target is about 2,000 tokens per second, from about 1,033 upstream. That's a goal, not a promise. Record every choice in NOTES.md.

Work through milestones M0-M8 in order. Each milestone has acceptance criteria. Don't start a milestone until the previous one passes, except that the page work in M7 can start on the Mac in parallel as soon as M2 produces a real trace. After each milestone, commit your work and write a short entry in NOTES.md with the results and numbers.

Hard constraints

• Secrets: Read RUNPOD_API_KEY and HF_TOKEN, and WANDB_API_KEY if you use it, from environment variables only. Never write a secret into any file, log, commit, or echoed command. Commit a .env.example that has placeholder values only, and add .env to .gitignore.
• Budget: The hard cap is $40 of RunPod spend. Every pod runs a watchdog that stops the pod after MAX_POD_HOURS hours. The default is 6.
• Compute: Run all CUDA builds, kernel runs, baselines, and quality evaluations only on RunPod. The local Mac runs trace analysis and replay, the web build, and the browser tests. Stop the pod when it's idle.
• Same-GPU rule: Every number in any comparison comes from the same GPU model, and ideally from the same pod session. Record the GPU, driver, CUDA version, clocks, power, and temperature with each run.
• Correctness gate: No speed number counts unless that exact build passed the correctness suite in the same session. For the bf16 port, that's a greedy token match against the reference, where the only allowed divergences are at positions where the reference's top-2 logit margin is below a tolerance that you record (bf16 accumulation order can flip near ties). For quantized variants, it's the quality gates. For every build, it includes determinism across 3 repeated runs. The timing harness lives outside the kernel. It times with CUDA events around the full decode loop, decode only with prefill excluded, after warm-up. The reason is KernelBench-Verified: with hidden tests, reported agent speedups collapsed, and documented reward hacks include a ReLU kernel that reported 374× and hardcoded shapes. Every speed number must come from a gated build, measured by a harness the kernel can't touch.
• Instrumented versus headline builds: X-ray compiles only behind a compile-time flag, and headline numbers come only from uninstrumented builds.
• Outward actions: Ask the user before you do any of the following: create the GitHub repository, push to the Hugging Face Hub, deploy to a VM, change DNS, or post anything publicly. Use the kashifulhaque GitHub account (gh auth switch -u kashifulhaque).
• Shared VM: vm.ifkash.dev runs other production apps behind one shared Caddy, which owns ports 80 and 443. Its config lives at ~/docs/caddy.
• The monolith container lives in ~/docs/monolith and must not publish any ports. It joins the external Docker network edge with a stable alias, such as monolith.
• Add the vhost only by appending to ~/docs/caddy/Caddyfile with >>. Never rewrite, rename, or replace that file. It's a single-file bind mount, and a rewrite orphans the inode, so caddy reload then reports "config is unchanged" while serving the old config.
• Docker on the VM has no BuildKit for plain docker build. Don't use COPY --chmod, Dockerfile heredocs, or RUN --mount. Multi-stage builds and COPY --from work.
• Don't stop, restart, reconfigure, or remove any other container, network, or vhost.
• Pod cleanup: When a pod isn't running a job, stop it. At the end of the project, terminate every pod that you created and report the total spend.
• Licenses: Keep upstream's MIT notice in every file that you copy or adapt, and credit AlpinDale's qwen_megakernel and MegaQwen in README.md and the page footer. Qwen3-0.6B is Apache-2.0; quantized weights keep that license and state that they're derived from it. llama.cpp (MIT) and vLLM (Apache-2.0) only run as baselines. State that monolith isn't affiliated with Qwen, NVIDIA, or AlpinDale.

Upstream facts

These facts come from AlpinDale's repo and blog post, the Qwen3-0.6B model card, and NVIDIA's RTX Blackwell whitepaper. Re-measure anything that you rely on.

• Upstream repo: https://github.com/AlpinDale/qwen_megakernel, MIT, created 2026-02-07 and last pushed 2026-02-25. Its description is "Aggressive decode optimizations for Qwen3-0.6B on RTX 5090." The top level holds csrc/, qwen_megakernel/, assets/, LICENSE, and requirements.txt. It builds on MegaQwen (https://github.com/Infatoshi/MegaQwen), which ran 527 tok/s on an RTX 3090.
• Upstream results (https://blog.alpindale.net/posts/5090_decode_optimization/): about 1,033 tok/s (0.97 ms per token) at batch size 1 with bf16 weights and no quantization. That's 71.2% of the 1,674 GB/s that the author measured on the RTX 5090. The initial MegaQwen port ran about 494 tok/s on the RTX 5090, then 813, 890, 905, and about 1,000 tok/s.
• Launch shape: 128 thread blocks of 512 threads each, on a GPU with 170 SMs.
• Phases and sync: Six phases per layer. In the author's words, the 71% ceiling "comes from the four remaining full-grid atomic barriers per layer (~2.2 us each x 4 = 8.8 us) plus two flag syncs (~0.2 us each)." A layer takes 28.0 µs measured, against a theoretical minimum of 18.8 µs, and "one-third of the per-layer time is synchronization overhead."
• Weights read per step: 1,192.1 MB, which is 880.9 MB for the 28 layers plus 311.2 MB for the LM head.
• Launch structure: The megakernel is launched as a regular (non-cooperative) kernel with custom atomic barriers. The author swept higher block counts and found 128 the sweet spot for the bf16 shapes. After the megakernel finishes, two small kernels compute the LM head argmax: 1,184 blocks of 256 threads, then a single 256-thread block for a tree reduction. The LM head plus startup is a fixed 216 µs of each 1,000 µs step.
• Tricks already in upstream: All 128 blocks compute the RMSNorm factor independently instead of waiting at a barrier. Two phases sync on monotonic flags instead of full barriers. "Productive spin": while 16 blocks compute attention, the other 112 issue prefetch.global.L2 for the next phases' weights (about 23 MB), which took the kernel from 905 to about 1,000 tok/s.
• Position scaling: Throughput stays above 970 tok/s out to position 200; the 96 MB L2 softens the KV cache growth.
• The author's next steps: fewer barriers ("the remaining four all need 128 blocks"), cheaper barriers, or quantization. This project does all three.
• Model (https://huggingface.co/Qwen/Qwen3-0.6B, Apache-2.0): hidden size 1024, intermediate size 3072, 28 layers, 16 query heads, 8 KV heads (GQA), head_dim 128, vocab 151,936, tied embeddings (the LM head is the embedding matrix), bf16, and max position 40,960. It uses QK-norm (RMSNorm on q and k per head) and RoPE.
• Weight counts (derived): Per layer, the attention projections hold about 6.29M weights and the MLP about 9.44M, so about 15.7M per layer and about 440M over 28 layers. The embedding and LM head hold about 155.6M. The total is about 596M.
• GPU (RTX 5090, GB202, sm_120): 170 SMs, 32 GB GDDR7 on a 512-bit bus, 1,792 GB/s spec, 96 MB L2, 2,407 MHz boost, 575 W. Per SM: a 256 KB register file, 128 KB unified L1 and shared memory, at most 99 KB shared memory per block (100 KB per SM), 1,536 threads, and 24 blocks.
• Tensor cores and ISA: Warp-level mma.sync only. There's no wgmma (sm_90a only), and no tcgen05 or TMEM (sm_100 family). FP4, FP6, and block-scaled mma.sync kinds (.kind::mxf4nvf4 and related kinds) need sm_120a or sm_120f. TMA exists; clusters and distributed shared memory exist; multicast isn't recommended on sm_120. setmaxnreg exists on sm_120a.
• Throughput: Dense tensor TFLOPS are 209.5 for FP16/BF16 with FP32 accumulate, 419 for FP8 with FP32 accumulate, and 1,676 for FP4; non-tensor FP32 is 104.8. Decode at batch 1 is matrix-vector work and is memory-bound, so tensor cores matter only for a batched speculative-verify stretch goal.
• Conversions: cvt to f16x2 from e4m3x2 exists since sm_89. Whether a native e2m1x2→f16x2 cvt exists on sm_120a is unverified; M0 checks it. The fallback is a 16-entry lookup through prmt or shared memory.
• Bigger sibling: The RTX PRO 6000 Blackwell has the same ISA with 188 SMs, 96 GB, 1,792 GB/s, and a 128 MB L2.

The byte budget and time model are as follows. The FP8 and NVFP4 lines are derived estimates:

bytes/token  bf16   ≈ 1.19 GB   (880.9 MB layers + 311.2 MB LM head, measured upstream)
             FP8    ≈ 0.60 GB   (per-output-channel scales, LM head FP8)
             NVFP4  ≈ 0.40 GB   (layers at 4.5 bits/weight, LM head FP8)
t_token      ≈ 28 × (bytes_layer / BW_eff + t_sync_layer) + bytes_head / BW_eff + t_argmax
upstream     28 × 28.0 µs, of which 8.8 µs is four grid barriers per layer

By this back-of-envelope estimate, at the measured 71% bandwidth and with the same four barriers, FP8 lands near 1,400 tok/s, because the barriers stay at 8.8 µs of every layer. Cutting sync to about 2 µs per layer takes FP8 to about 1,900 tok/s and NVFP4 to about 2,500 or more. That's why the milestones run FP8, then barriers, then NVFP4.

Design decisions

• Fork, don't rewrite: Fork upstream at a pinned commit, and keep its phase structure through M3. Restructure only in M4. Record the commit hash in NOTES.md and THIRD_PARTY.md.
• FP8 format: E4M3 weights with one fp32 (or bf16) scale per output row, symmetric, round to nearest. Activations stay bf16 or fp32 (W8A16). Dequantize in registers: load 16 bytes, cvt e4m3x2→f16x2 (sm_89 and later), apply the scale, and FMA in fp32. The embedding lookup reads the FP8 row and dequantizes it, because the matrix is tied and there's one copy. Record the scale dtype that you pick.
• NVFP4 format: E2M1 values, one E4M3 scale per 16 weights, and one fp32 scale per tensor. Quantize with round to nearest plus a per-block scale search that minimizes MSE. Optionally, compare against a calibrated checkpoint from NVIDIA ModelOpt or llm-compressor if one runs in reasonable time; whether it does is unverified, so record what happens. Mixed precision is allowed: a per-layer sensitivity sweep decides which matrices stay FP8. The LM head stays FP8. Record the final per-matrix format map.
• Weight layout: Pack offline into the exact order in which each block streams its tiles: row-tiled, 16-byte aligned, with scales interleaved per tile, so every warp issues coalesced 128-bit loads. Record the layout in NOTES.md.
• KV cache: Stays bf16. Cap the context at a length that you pick, at least 2,048, and record it.
• Sync experiments: In M4, explore these in order, and record each experiment's tok/s and X-ray trace: 1. Phase fusion that removes another barrier by recomputing a cheap result per block. Upstream already does this for the RMSNorm factor, so use the X-ray to find the next candidate. 2. Hazy-style counters: each producer tile increments a counter, and each consumer waits only on the tiles that it reads, with release/acquire ordering (st.release and ld.acquire, or atom with .release and .acquire, at .gpu scope). 3. Launch shape: re-sweep the block count, up to all 170 SMs. Upstream found 128 best for bf16, but quantization changes the balance between bandwidth and barrier cost. 4. More productive spin: upstream prefetches into L2 only during attention; extend it to the other waits. 5. Fold the LM head and argmax into the megakernel, so that one launch produces each token. In FP8 the LM head reads about 0.16 GB, so it's also a large phase to schedule well.
• Timing: Use %globaltimer (nanoseconds, but its update granularity might be coarse, so measure it in M0) plus clock64 and %smid. Lane 0 of warp 0 writes one stamp pair per block per phase to a preallocated device buffer. Align the per-SM clocks offline. Record the alignment method.
• Benchmark protocol: A fixed set of 20 prompts, prefill excluded, 256 generated tokens, greedy decoding. Report tok/s over positions 1-256, plus a position sweep at 16, 128, 512, 1024, and 2048. Report the median of 5 runs with the minimum and maximum, plus effective bandwidth, which is bytes read divided by time. Commit the prompt set.
• Quality protocol: WikiText-2 test perplexity with a fixed stride and context (record both); greedy agreement on the first 64 tokens over 100 prompts against bf16; and mean KL divergence of the next-token distribution over 10k positions. The gates are as follows, and you record the final thresholds:
• FP8 perplexity is within +2% of bf16, and first-64-token agreement is 90% or more.
• NVFP4 is reported as is, and labeled experimental if perplexity rises more than 10%.
• Speed and cost (estimates): Predict each variant with the byte-budget model, and record the prediction next to the measurement. The project is about 15-25 pod-hours, about $15-25.

Tech stack

Pin every version. The stack is as follows:

• Kernel: CUDA 13.x toolkit (pin it; CUDA Toolkit 13.4 Update 1 with PTX ISA 9.4 is the reference point), nvcc -arch=sm_120a, and C++17.
• Python: PyTorch (pinned, CUDA 13 build) for the extension and the reference path, transformers for the Hugging Face baseline and reference logits, datasets for WikiText-2, numpy, and matplotlib, managed with uv.
• Baselines: llama.cpp at a pinned commit, built with CUDA for sm_120, and vLLM (pinned). llama.cpp merged native Blackwell NVFP4 support in April 2026 according to a secondhand source (https://insiderllm.com/guides/fp4-inference-llamacpp-nvfp4-mxfp4/); M1 confirms whether it runs Qwen3-0.6B on sm_120.
• Profiling: nsys for timelines, ncu only if M0 shows that counters work, and compute-sanitizer for racecheck and memcheck on reduced configs.
• Web: TypeScript bundled with vite, Canvas 2D (or WebGL if trace size needs it), and no framework required. The page reads gzipped JSON traces.
• Browser tests: Playwright with Chromium.
• Pods: the runpod Python SDK, with runpodctl inside pods.

Repository layout

Create the following layout:

monolith/
  pyproject.toml  .env.example  .gitignore  README.md  NOTES.md  LICENSE  THIRD_PARTY.md
  infra/pod.py  infra/watchdog.sh  infra/bootstrap.sh
  probes/             # M0: cvt_e2m1.cu, timer_res.cu, bandwidth.cu, ncu_check.sh
  upstream/           # pinned fork of qwen_megakernel (MIT notice kept)
  kernel/
    megakernel.cu     # phases, barriers/counters, dequant GEMV, attention, LM head argmax
    dequant.cuh       # fp8 and nvfp4 register dequant helpers
    sync.cuh          # grid barrier, counters, release/acquire helpers
    xray.cuh          # MONOLITH_XRAY stamps: globaltimer, clock64, smid
    bindings.cpp      # torch extension
  quant/pack_fp8.py  quant/pack_nvfp4.py  quant/sensitivity.py
  bench/decode.py     # timing harness (CUDA events, decode only)
  bench/baselines/    # hf.py, llamacpp.sh, vllm.py
  eval/correctness.py  eval/quality.py   # token match, ppl, agreement, KL
  xray/analyze.py     # per-phase time, barrier wait, idle per SM, bandwidth
  xray/export.py      # compact trace JSON for the page
  web/src/            # race.ts, xray.ts (Gantt), budget.ts, table.ts, tour.ts, main.ts
  web/public/traces/  web/tests/
  deploy/Dockerfile  deploy/compose.yml  deploy/nginx.conf
  results/  post/draft.md

M0: Infrastructure and probes

Tasks:

1. Write infra/pod.py. It uses the runpod SDK and reads RUNPOD_API_KEY from the environment. It supports the following subcommands: - create: Creates a pod. The GPU preference order is RTX 5090, then RTX PRO 6000. Resolve the GPU type IDs at run time by querying the available GPU types. Use an official RunPod image with CUDA 12.8 or later, on a host whose driver is new enough for the pinned toolkit. GeForce cards get no CUDA forward compatibility, so filter by driver: R580 or later for CUDA 13.x. Set a 50 GB volume at /workspace and expose SSH. - status, stop, and terminate. - ssh-info: Prints the SSH command. - cost: Prints the uptime and spend for every pod whose name has the prefix monolith-.
2. Write infra/watchdog.sh. It sleeps for MAX_POD_HOURS, then runs runpodctl stop pod $RUNPOD_POD_ID.
3. Write infra/bootstrap.sh. It clones the repo, installs the pinned CUDA toolkit if the image lacks it, runs uv sync, and starts the watchdog.
4. Write the probes in probes/ and run each one on the pod: - ncu_check.sh: Runs ncu --section SpeedOfLight on a trivial kernel. RunPod pods are unprivileged containers, and RunPod Discord threads report ERR_NVGPUCTRPERM for hardware counters and say that --cap-add=SYS_ADMIN isn't allowed. That's unverified for any specific host, so record what this host does. If counters are blocked, nsys and in-kernel timers still work. - timer_res.cu: Measures the update granularity of %globaltimer and the rate of clock64, and checks how consistent clock64 is across SMs. - cvt_e2m1.cu: Compiles with -arch=sm_120a and runs cvt e4m3x2→f16x2 and e2m1x2→f16x2 against a CPU reference for every input value. If the e2m1 form doesn't compile or doesn't match, record that M5 uses the 16-entry lookup through prmt or shared memory. - bandwidth.cu: A streaming-read kernel that measures achievable DRAM bandwidth. Compare it with upstream's 1,674 GB/s and the 1,792 GB/s spec, and use the result as BW_eff in the byte-budget model.

Acceptance criteria:

• A pod comes up, bootstrap.sh finishes without errors, and nvidia-smi shows the GPU and the driver version.
• A test run with MAX_POD_HOURS=0.05 stops the pod within 5 minutes.
• NOTES.md records every probe result: ncu status, timer granularity, clock64 rate, both cvt results, and measured bandwidth.

M1: Reproduce upstream and baselines

Tasks:

• Clone upstream into upstream/ at a pinned commit, build it, and record the hash.
• Write eval/correctness.py. It runs the reference in transformers with greedy decoding and compares tokens against the kernel for the 20-prompt set at 256 tokens each. It flags every divergence and records the reference's top-2 logit margin there.
• Write bench/decode.py with the benchmark protocol from the design decisions. It times with CUDA events around the full decode loop, after warm-up, with prefill excluded.
• Write the baselines in bench/baselines/:
• hf.py: Hugging Face transformers in bf16 with a static cache, and with torch.compile if it works. Record whether it did.
• llamacpp.sh: llama.cpp in BF16 and Q8_0, and NVFP4 if it runs Qwen3-0.6B on sm_120. If NVFP4 doesn't run, drop that row and record why.
• vllm.py: vLLM in bf16 with CUDA graphs at batch 1.

Acceptance criteria:

• Upstream reaches within 5% of about 1,033 tok/s, or NOTES.md explains the deviation.
• Upstream matches Hugging Face greedy tokens on 20 prompts × 256 tokens, except at near-tie positions below the recorded margin tolerance, and NOTES.md reports the match rate.
• results/m1_baselines.json holds the baseline table, measured with the same protocol on the same pod.

M2: X-ray

Tasks:

• Write kernel/xray.cuh. When MONOLITH_XRAY is defined, lane 0 of warp 0 in each block writes a start and end stamp (%globaltimer, clock64, and %smid) for every phase to a preallocated device buffer. When it isn't defined, the macros compile to nothing.
• Add the host-side dump: after a decode run, copy the trace buffer for one chosen token position to the host.
• Write xray/analyze.py. It aligns per-SM clocks, then reports per-phase time, barrier wait per block, idle time per SM, and effective bandwidth per phase. Add a quick matplotlib Gantt chart for checking traces without the page.
• Write xray/export.py. It writes a compact trace JSON for the page, gzipped.

Acceptance criteria:

• The instrumented build's tok/s is within 3% of the uninstrumented build's.
• The summed phase time matches the measured per-token time within 5%.
• The analyzer reproduces upstream's finding. NOTES.md reports the barrier share of layer time next to upstream's "one-third."
• A real trace JSON is committed for M7.

M3: FP8 weights

Tasks:

• Write quant/pack_fp8.py. It quantizes every projection and the tied embedding and LM head to E4M3 with per-output-row scales, and packs them in the streaming layout from the design decisions.
• Write kernel/dequant.cuh with the FP8 register dequant helper, and switch every matrix-vector product and the LM head to it. The embedding lookup dequantizes the FP8 row.
• Write eval/quality.py with the quality protocol: WikiText-2 perplexity, first-64-token greedy agreement over 100 prompts, and mean KL divergence over 10k positions, all against bf16.
• If argmax agreement drops below the gate, keep the LM head in bf16 (about 0.16 GB of extra reads per token), and record it.
• Record the X-ray trace of the FP8 kernel, and compare the measured tok/s with the byte-budget prediction.

Acceptance criteria:

• FP8 passes the quality gates.
• FP8 tok/s beats bf16. The estimate is about 1,300-1,400 tok/s.
• The FP8 X-ray trace shows the barrier share of layer time rising compared with bf16, and NOTES.md reports both shares.

M4: Fewer barriers

Tasks:

• Write kernel/sync.cuh with the grid barrier, per-tile counters, and release/acquire helpers.
• Run the sync experiments from the design decisions one at a time, in order. For each one, run the correctness gate, a determinism test over 3 repeated runs, compute-sanitizer racecheck on a reduced config, the benchmark, and an X-ray trace.
• Keep an experiment only if it passes every check and improves tok/s. Otherwise revert it, and record why.

Acceptance criteria:

• The FP8 kernel reaches at least 1,700 tok/s with all gates passing. The goal is about 2,000.
• The LM head and argmax run inside the megakernel, so one launch produces each token, or NOTES.md shows the measurement that made you keep them separate.
• NOTES.md has a table of every experiment with its tok/s, barrier share, and whether you kept or reverted it.

M5: NVFP4

If time runs out, skip this milestone and say so in NOTES.md and the final report. FP8 is the headline.

Tasks:

• Write quant/pack_nvfp4.py: E2M1 values, one E4M3 scale per 16 weights, and one fp32 scale per tensor, with round to nearest and the per-block MSE scale search. Optionally compare against a ModelOpt or llm-compressor checkpoint, and record whether it ran.
• Write quant/sensitivity.py. It quantizes one matrix at a time to NVFP4 and measures the quality change, then proposes a mixed-precision config where sensitive matrices stay FP8.
• Add the e2m1 dequant path to kernel/dequant.cuh, using the native cvt or the 16-entry lookup, based on the M0 probe result.
• Keep the LM head in FP8.

Acceptance criteria:

• The NVFP4 kernel runs correctly under the correctness gate.
• Quality is measured and reported honestly. If perplexity rises more than 10% over bf16, the variant is labeled experimental everywhere.
• NVFP4 tok/s is recorded next to its byte-budget prediction.

M6: Final benchmark matrix

Tasks:

• In one pod session, run every variant (upstream bf16, FP8, FP8 with fewer barriers, and NVFP4 if it exists) and every baseline with the full benchmark protocol, including the position sweep.
• Compute tokens per joule from nvidia-smi power readings, and report bytes per token and effective bandwidth for every monolith variant.
• Export one representative token trace per variant at position 128, from the instrumented build, and export per-token timestamps from the uninstrumented builds and baselines for the race.

Acceptance criteria:

• results/final.json and results/final.md hold the full matrix with quality columns.
• The exported traces total 5 MB or less, gzipped.

M7: The demo page

The page can start on the Mac as soon as M2 commits a real trace.

Tasks:

• Race: Four token streams, from monolith and the three baselines, replay at their recorded per-token timings, with a tok/s counter on each.
• X-ray: An SM × time Gantt chart of one token, with one row per block, one color per phase, and hatched barrier waits. Support zoom and pan, and show phase, SM, and duration when the pointer is over a bar. A version switcher steps through bf16 upstream, FP8, FP8 with fewer barriers, and NVFP4.
• Budget bar: For each version, split one token into weight reads, barrier waits, attention, and the LM head, next to the roofline minimum from the byte-budget model.
• Results table: Speed, bytes per token, effective bandwidth, and tokens per joule, with perplexity, agreement, and KL divergence columns next to them. Mark experimental variants.
• Code tour: The dequant inner loop and the counter wait, with short annotations.
• Methodology notes: The same-GPU rule, the correctness gate, and the benchmark protocol.
• Footer: Credits for AlpinDale's qwen_megakernel and MegaQwen, the Qwen3-0.6B license, and "not affiliated with Qwen, NVIDIA, or AlpinDale."
• Mobile: Stack the panels, and support touch pan and zoom on the Gantt chart.

Acceptance criteria:

• A Playwright smoke test loads the page, plays the race, switches X-ray versions, and holds the pointer over a block to show its details.
• The page works in Chrome and Safari on macOS.
• The page loads in under 2 seconds on a normal connection.

M8: Deploy and write-up

Tasks:

• Write deploy/Dockerfile: nginx:alpine serving web/dist, buildable with the legacy builder. Serve the precompressed traces with gzip_static on.
• Write deploy/compose.yml. It publishes no ports and joins the external network edge with the alias monolith.
• Ask the user before deploying. The user confirms a domain such as monolith.ifkash.dev and adds a non-proxied Cloudflare A record.
• After the user approves, copy the build and deploy/ to ~/docs/monolith, and run docker compose up -d there. Then compare the live and host Caddyfiles with docker exec caddy cat /etc/caddy/Caddyfile | diff - ~/docs/caddy/Caddyfile. If they differ, stop and ask the user. Otherwise, back up, append, validate, and reload:

cd ~/docs/caddy cp Caddyfile "Caddyfile.bak-monolith-$(date +%Y%m%d)" printf '\nmonolith.ifkash.dev {\n\treverse_proxy monolith:80\n}\n' >> Caddyfile docker exec -i caddy caddy validate --config - --adapter caddyfile < Caddyfile docker exec -i caddy caddy reload --config - --adapter caddyfile < Caddyfile

Omit any tls block: for a non-proxied A record, the default ACME HTTP-01 challenge works. - Write README.md. It explains what the project is, how to reproduce each milestone with one command per milestone, the results tables, every deviation from this plan, and the credits and license notes. - Write post/draft.md, a blog post of 1,200-1,800 words. Structure it as follows: - The hook: one kernel, one token, 170 SMs. - What a megakernel is, in plain words. - The X-ray, and the barrier twist it revealed. - FP8 and NVFP4 dequantized in registers. - Killing barriers: fusion, counters, launch shape, and one launch per token. - Honest benchmarks: the same GPU, the correctness gate, and the baselines. - Limitations: batch 1 only, one model, one GPU, and weight-only quantization. - Ask the user before you push anything. After the user approves, push the repo to weights-and-wires/monolith and upload the packed weights to the Hugging Face Hub under weights-and-wires/monolith-qwen3-0.6b, with the Apache-2.0 license and a note that they're derived from Qwen3-0.6B. - Terminate every pod, and write the final spend to NOTES.md. The owner lists the project on projects.dotslasha.me after the deploy; that isn't your job.

Acceptance criteria:

• The deployed URL loads over HTTPS and plays the race and the X-ray in Chrome.
• Every other site on the VM still responds as it did before the deploy.
• No pods are left running, and the total spend is less than $40.

Final report to the user

When you finish, report the following:

• The probe results from M0: ncu status, timer granularity, both cvt results, and measured bandwidth.
• The upstream reproduction and the baseline numbers, from M1.
• The X-ray overhead and the barrier share of layer time next to upstream's one-third, from M2.
• The FP8 tok/s and quality numbers, from M3.
• The M4 experiment table and the best FP8 tok/s.
• The NVFP4 result, or a note that you skipped M5.
• The final benchmark matrix, from M6.
• The URL, if the site is deployed.
• The total spend.
• Anything that you skipped or that failed.