Recent content on gau-nernst's blog
Some time ago, I explored using CuteDSL to write some kernels for vLLM to enjoy the JIT benefits (ability to specialize on more parameters, don’t need to recompile vLLM from source during development, etc…). I have always known that we can write CUDA C++ way (SIMT style) in CuteDSL: there are cute.arch.thread_idx() / cute.arch.block_idx() (equivalent of threadIdx / blockIdx in CUDA…
tcgen05 is the set of PTX instructions to program Tensor Cores on the latest NVIDIA Blackwell GPUs ( sm100 , not to be confused with consumer Blackwell sm120 ). At the time of writing, I couldn’t find a Blackwell tutorial in plain CUDA C++ with PTX, even though such exist for Ampere ( alexarmbr’s and spatters’ ) and Hopper ( Pranjal’s ). So let’s write one,…
Last month, I participated in the AMD Distributed Challenge, hosted by GPU MODE . This was very exciting for me as it was the first time I learned how to write a multi-GPU kernel! Although I had a brief understanding of how DDP and FSDP worked under the hood via collective primitives like all-reduce and reduce-scatter, I didn’t know it was possible to perform remote memory access directly…
Recently I tweeted about realistic Speed-of-Light (SOL) of 5090 and RTX PRO 6000 for some dtypes, and mobicham asked me about FP8 MMA with FP16 accumulation . I of last year would turn to Triton for this - it’s trivial to change the accumulation dtype of tl.dot() . However, I roughly know how to write a fast matmul kernel now, so why not do it myself! In addition, I have been tinkering…
In this post, I will walkthrough how I learned to implement Flash Attention for 5090 in CUDA C++. The main objective is to learn writing attention in CUDA C++, since many features are not available in Triton , such as MXFP8 / NVFP4 MMA for sm120. I also feel this is a natural next step after learning about matmul kernels. Lastly, there are many excellent blogposts on writing fast matmul kernels,…