Raj (🦀/gpu)

currently into gpu programming and ml systems
rust | gsoc'25
[email protected]

[Raj (🦀/gpu) @Raj_ghj]

@funinkina It saves time for me and I feel fast , so whtvr

[Raj (🦀/gpu) @Raj_ghj]

@Shreyans_23 How much ?

[Soham Parekh - in/acc @realsohamparekh]

Recommend Wispr Flow to a friend. Got hit with, “I’m not menstruating.”

Incredible product.
Unfortunate name.

[Raj (🦀/gpu) @Raj_ghj]

💅💅

[Sam Altman @sama]

there are a lot of benchmarks that suggest 5.6 sol is the best model in the world right now, but the most reliable way to tell is that elon is obsessed with me again

[Samay @Samaytwt]

This is how Apple search engine works btw.

[Raj (🦀/gpu) @Raj_ghj]

@ruskaruma08 Is "not interested" not working

[Raj (🦀/gpu) @Raj_ghj]

@SwastKumar Linux muslim 😭

[Raj (🦀/gpu) @Raj_ghj]

Homecoming - mr west

[luffy @0xluffy]

fable 5 reading the code written by opus 4.8 after going afk for 19 days

[Raj (🦀/gpu) @Raj_ghj]

@levidiamode you still didnt answer my question from the last post

[Raj (🦀/gpu) @Raj_ghj]

@ID_AA_Carmack @ptrCast maybe you should test it ... You said ai is performing better in low level code... Might go even lower maybe

[Raj (🦀/gpu) @Raj_ghj]

@ishaans08 Why ??

[Raj (🦀/gpu) @Raj_ghj]

The guy wrote something so simple yet relatable and effective, it's as if I'm reading my mind

Here are some of the things that I love about this write-up

[Raj (🦀/gpu) @Raj_ghj]

@levidiamode did you do it all by yourself or was it agentc

[Raj (🦀/gpu) @Raj_ghj]

If a person speaks a lot , there is more chance that they speak 90% bullshit with 10% interesting stuff they don't even know deeply.

[Raj (🦀/gpu) @Raj_ghj]

Keep your god complex to yourself
This world has seen enough people like you.

[Zhihu Frontier @ZhihuFrontier]

⚡ GPU programming changed because Tensor Cores became too fast to feed
Zhihu contributor THU-PACMAN实验室 shared a sharp breakdown of how NVIDIA GPU programming moved from Volta to Blackwell.
The headline is not “FP16 became FP4” or “TFLOPS went up.” The real shift is deeper: a CUDA kernel is no longer just a group of threads running in sync. It is becoming an asynchronous dataflow program across compute units, memory engines, barriers, buffers, and layouts.

⚙ The bottleneck moved to feeding Tensor Cores
From V100 to B100, peak FP16 throughput exploded:
V100: 112 TFLOPS, 0.90 TB/s HBM, ~124 FLOPS/Byte
A100: 312 TFLOPS, 2.04 TB/s HBM, ~153 FLOPS/Byte
H100: 990 TFLOPS, 3.35 TB/s HBM, ~296 FLOPS/Byte
B100: ~1800 TFLOPS, 8.00 TB/s HBM, ~225 FLOPS/Byte
On Hopper, even if HBM is fully saturated, each byte from global memory has to support roughly 300 floating-point operations.
So the hard part is no longer “do we have matrix units?” It is how to keep them continuously fed.

🔄 Data movement became a hardware pipeline
Volta still looked like classic CUDA: threads computed addresses, loaded from global memory, staged through registers, wrote into shared memory, then synchronized.
Ampere added cp.async, letting data move from global memory to shared memory without going through registers. Hopper pushed the model further with TMA, where hardware handles tile-level address generation, strides, boundaries, and background movement.
Blackwell adds TMEM, a dedicated tensor memory layer, and makes the kernel look even less like a normal thread program.

🧠 Matrix instructions also changed their meaning
Volta wmma: a warp synchronously computes together.
Ampere mma.sync: more control over shape and dtype, with explicit shared-memory layouts.
Hopper https://t.co/4qAg59kfTu_async: Tensor Cores become an asynchronous compute agent.
Blackwell https://t.co/LmZqumopS2: accumulation can live in TMEM, a new 256KB/SM tensor storage layer.
The mental model changes from “all threads participate in this computation” to “launch work to a hardware unit, track its completion, and manage where the result lives.”

🚦 Synchronization is no longer just “wait for threads”
Traditional CUDA synchronization was mostly control-flow synchronization: __syncthreads() means everyone in the block has arrived.
But with TMA, WGMMA, and Blackwell async MMA, the important question is different: who produced this data, who consumes it, which async proxy owns the operation, what signal proves the data is ready, and when the buffer can be reused.
In modern kernels, synchronization is becoming fine-grained data dependency management.

📉 Low precision is not just changing dtype
FP8 and FP4 are often described as smaller types. But on Blackwell, low precision becomes a whole constraint system.
For block-scaled MMA, the compiler must jointly reason about dtype, scale vector size, scale tensor layout, operand packing, alignment, physical layout, accumulator type, tile shape, dispatch policy, and which tcgen05 instruction is legal.
So instruction selection is no longer a late backend peephole. It is tied to dtype, packing, scale placement, layout, accumulator choice, and epilogue design.

🚀 FlashAttention shows the new programming model clearly
FlashAttention-3 on H100 treats the GPU like a set of schedulable hardware units: one warpgroup runs WGMMA on Tensor Cores, another overlaps softmax on CUDA Cores, and TMA prefetches the next tile in the background.
This “ping-pong scheduling” raised utilization sharply: FlashAttention-2 used about 35% of H100 peak, FA3 FP16 reached 740 TFLOPS, and FA3 FP8 approached 1.2 PFLOPS.
FlashMLA shows the same lesson from another angle. MLA decoding can require 32,768 32-bit registers for one 64×512 output tile, while an SM has only 65,536 registers. DeepSeek’s “seesaw scheduling” splits the output, alternates warpgroups, overlaps Tensor Core and CUDA Core work, and starts GEMM as soon as the first fine-grained TMA copy arrives.
👉 Modern kernel optimization is no longer just choosing an algorithm. It is designing a schedule for hardware units.

🧩 What this means for compilers
A compiler or DSL for modern AI chips needs to describe at least five things: tile shape, memory movement, physical layout, async producer-consumer dependencies, and the schedule of copy, compute, sync, and writeback.
Without these, correctness is hard to verify and performance is hard to explain.
This also means a unified frontend does not magically produce optimal backend code. Different hardware still needs specialized schedules, synchronization strategies, layouts, and resource models.
The biggest change from Volta to Blackwell is not a faster instruction. It is a new way to think about a GPU kernel: from synchronous thread cooperation to asynchronous dataflow across hardware components.

🔗 Full analysis:
https://t.co/RoKE6Runok

#GPU #CUDA #AIInfra #Compiler #HPC #Blackwell #FlashAttention #Triton

[Raj (🦀/gpu) @Raj_ghj]

@SwastKumar @pepsi Whatttt

[Raj (🦀/gpu) @Raj_ghj]

"I despise the idea of borrowing frameworks and mental models. Because a borrowed identity is a comfortable cage."

[Raj (🦀/gpu) @Raj_ghj]

A language that doesn't affect the way you think about programming is not worth knowing

– Alan perlis (first ACM Turing Award winner)