Writing
Notes from systems work, experiments worth preserving, and the occasional paper.
Articles
The Training and RL Infrastructure Behind MAI-Thinking-1Hardware, numerics, the YOLO training framework and the Rocket RL stack behind Microsoft AI's 35B-active, 1T-total MoE, read from its technical report.October 13, 2026Porting Block-Scaled GEMM Kernels from Blackwell to RubinWhat changes for a tcgen05 GEMM on sm_107a, read from FlashInfer's Rubin kernels: wider K, packed FP4, bigger SMEM and TMEM, new scale formats, and a mixed-cluster scheduler.October 13, 2026How Tokens Reach Their Experts During LLM InferenceHow SGLang moves each token to the GPUs that hold its experts and back: the twelve values of --moe-a2a-backend, from a plain all-gather to DeepEP and FlashInfer, and when each one pays off.October 11, 2026How GPUs Divide Up the Requests During LLM InferenceOn GLM-5.3, DP attention lets 4 GPUs hold 3.3 times more context, and makes a lone request almost twice as slow. Where that time goes, and at which load the trade is worth it.October 11, 2026Serving Kimi K3 End to End: 8 GPUs, 32 GPUs, and a RackServing Kimi K3 with SGLang at three scales: one 8-GPU server measured, 32 GPUs as one wide expert-parallel domain, and an NVL72 rack with prefill and decode split.October 11, 2026How GPUs Fuse Communication and Computation During LLM InferenceDeepGEMM's Mega MoE runs dispatch, both expert GEMMs, and combine as one kernel per GPU. Where the saved time comes from, and what its multi-gigabyte buffer costs.October 11, 2026What Matters in a ModelOpt Quantization RecipeHow FP8 and NVFP4 work, and what each field of a ModelOpt quantization recipe changed on a model served with SGLang.October 11, 2026How a Row Meets Its ExpertsMixture-of-Experts parallelism in SGLang, from first principles: five ways for a row to reach its experts across GPUs, measured for decoding and prefill on 8 GPUs.October 11, 2026How an MoE Model Picks Its Experts During LLM InferenceThe small matrix product and top-k that route every token in a Mixture-of-Experts layer: what it costs in SGLang, and how far a fused kernel takes it.October 11, 2026How GPUs Avoid Waiting for the CPU During LLM InferenceSGLang's overlap scheduler, from first principles: how the next step is queued before the CPU has seen the current one, and what speculative decoding adds.October 11, 2026How GPUs Share a Long Context During LLM InferencePrefill and decode context parallelism in SGLang, from first principles: why splitting a long context by token takes two separate features, and what each one costs.October 8, 2026Keeping Speculative Decoding Fast During RL TrainingWhy a drafter goes stale while RL changes the model it drafts for, and how verl-SpeCo, slime, Miles, and SpecForge keep it up to date.October 2, 2026