Have you ever been so excited about the thing you're working on that you're unable to focus on actually working on it?

The question is not who's Bob, the question is where's the reverse-engineered floorplan of the H100 chip

TIL: sometimes you have to commit programming crimes to please the compiler gods

If you ever implemented swizzling you know how helpful this is

Great analysis of embedding models by Piotr Mazurek

Penny is now a working group! If you want to make the world a better place by creating a well documented, performant and minimalistic AllReduce example join the GPU MODE discord server!

New in-depth blog post time: "Inside NVIDIA GPUs: Anatomy of high performance matmul kernels". If you want to deeply understand how one writes state of the art matmul kernels in CUDA read along. (Remember matmul is the single most important operation that transformers execute

Played around a bit with oneshot allreduce, already getting good results on small buffers(80% of NCCL) and it's just a lazy version, should be able to optimize this further

Interesting shift in GPU programming is the shift from parallel to parallel + async. Ampere was async loads Hopper was async loads + async wgmma ops Blackwell doesn't return values to registers but tensor memory When do we shrink the register file to save chip space?

Did some work on speeding up oneshot reduction in Penny, huge gains on small buffers. Time to crack midsize buffers and update the worklog

Visualisation of achieved MFU for FP16 matmul across different shapes, you can clearly see the effect of Wave Quantization and Tile Quantization AKA pick your matrix shape wisely

Fact: 90% of researcheres kill their run before the model is about to recover