WingEdge777

[CUDA in Practice] Hand-Rolled Flash Decoding on SM120: Beating flashinfer.single_decode_with_kv_cache

[CUDA in Practice] Hand-Rolled Flash Decoding on SM120: Beating flashinfer.single_decode_with_kv_cache

codeMay 21, 2026

Hand-rolled SM120 flash decoding kernel for single-query, long KV-cache decode attention — optimized to beat flashinfer's single_decode_with_kv_cache.

[CUDA in Practice] FMHA on SM120: Beating torch.sdpa (FlashAttention-2)

[CUDA in Practice] FMHA on SM120: Beating torch.sdpa (FlashAttention-2)

codeMay 19, 2026

Hand-crafted FMHA on SM120 with TMA, ldmatrix, and mma — a prefill Flash Attention kernel that beats torch.sdpa and FlashAttention-2 on RTX 5060.

Distributed Inference with PyTorch from First Principles: DP, TP, and PP in Less Than 200 Lines

Distributed Inference with PyTorch from First Principles: DP, TP, and PP in Less Than 200 Lines

codeMay 15, 2026

Rebuild LLM distributed inference from scratch with PyTorch: Data Parallelism, Tensor Parallelism, and Pipeline Parallelism in under 200 lines on 2 GPUs — without Megatron or vLLM.

[CUDA in Practice] HGEMM SM120 — Micro-Sculpture Warfare in 100KB SMEM: Tensor Core, TMA, ldmatrix, mma

[CUDA in Practice] HGEMM SM120 — Micro-Sculpture Warfare in 100KB SMEM: Tensor Core, TMA, ldmatrix, mma

codeMay 10, 2026

HGEMM on RTX 5060 SM120: squeezing peak performance from 100KB shared memory with TMA, ldmatrix, mma, and Tensor Core tuning.

[CUDA in Practice] HGEMM — Beating cuBLAS: Tensor Core, cp.async, ldmatrix, mma

[CUDA in Practice] HGEMM — Beating cuBLAS: Tensor Core, cp.async, ldmatrix, mma

codeMay 10, 2026

FP16/BF16 HGEMM hand-tuned with cp.async, ldmatrix, mma, and swizzle — a full walkthrough of beating cuBLAS on RTX 5060 for half-precision GEMM.

[CUDA in Practice] SGEMM TF32 — Beating cuBLAS with Tensor Cores, cp.async, ldmatrix & mma

[CUDA in Practice] SGEMM TF32 — Beating cuBLAS with Tensor Cores, cp.async, ldmatrix & mma

codeMay 9, 2026

TF32 Tensor Core SGEMM deep dive: cp.async, ldmatrix, mma, and XOR swizzle derivations — hand-tuned to challenge cuBLAS at specific problem sizes.

[CUDA Basics] Understanding CUDA's "Nonexistent" Memory Tier: Local Memory

[CUDA Basics] Understanding CUDA's "Nonexistent" Memory Tier: Local Memory

codeApril 1, 2026

CUDA primer on local memory: logically per-thread but physically in global memory, when the compiler spills to it, and how to avoid the performance hit.

[CUDA in Practice] Safe Online Softmax — A Must-Know for Interviews: Arbitrary hidden_size, One/Two Pass, Trade-offs, Split-K

[CUDA in Practice] Safe Online Softmax — A Must-Know for Interviews: Arbitrary hidden_size, One/Two Pass, Trade-offs, Split-K

codeMarch 31, 2026

A CUDA guide to safe online softmax across arbitrary hidden_size, one-pass vs two-pass, split-k, and the trade-offs behind each — a common interview topic.

[CUDA in Practice] SGEMM — Beating cuBLAS: A Deep Dive into Peak-Performance Matrix Multiplication in Pure CUDA C++

[CUDA in Practice] SGEMM — Beating cuBLAS: A Deep Dive into Peak-Performance Matrix Multiplication in Pure CUDA C++

codeMarch 5, 2026

Hand-crafted FP32 SGEMM in pure CUDA C++: from naive tiling to double-buffering and swizzle, beating cuBLAS on RTX 5060 with a full optimization walkthrough.

[CUDA in Practice] Matrix Transpose — From Padding to XOR Swizzle: The Art of Shared Memory Optimization

[CUDA in Practice] Matrix Transpose — From Padding to XOR Swizzle: The Art of Shared Memory Optimization

codeFebruary 13, 2026

CUDA matrix transpose optimization: from coalesced vs strided access through shared-memory tiling, padding, and XOR swizzle to eliminate bank conflicts.

1 2