I was surprised to see 5090's theoretical BF16 TFLOPs at just 209.5. That's not even 10% of the server Blackwell (B200 is 2250, and GB200 is 2500). B200 costs around $30-40k per GPU, so they are pretty close in performance per dollar.
Starting with 4090, NVIDIA limits the performance of tensor cores on gaming cards, specifically for ops that might be used in ML training. FP8 and FP16 matmuls run at full speed if accumulating in FP16 (I've never seen anyone use this), but only half speed when accumulating in FP32. This restriction is not present for lower precision matmuls like FP4, and is removed entirely on the workstation-class cards like RTX Pro 6000.
It doesn't seem worth it to use NVIDIA gaming cards as a "cheaper FLOPs" alternative anymore (e.g. diffusion models could have been cheaper to run on 3090 than A100). They are generous with memory bandwidth though, nearly 2TB/s on 5090 is amazing!
Is there really that big a different in TFLOPS between the GB100 and GB202 chips?
The GB100 has fewer SMs than the GB202, so I'm confused about where the 10x performance would be coming from?
You're asking a really good question but it's not a question with an easy answer.
There's a lot more to performance computing than FLOPs. FLOPs are you good high level easy to understand metric but it's a small part of the story when you're in the weeds.
To help make sense of this, look at CPU frequencies. I think most people on HN know that two CPU with the same frequency can have dramatically different outcomes on benchmarks, right? You might know how some of these come down to things like IPC (instructions per cycle) or the cache structures. There's even more but we know it's not so easy to measure, right?
On a GPU all that is true but there's only more complexity. Your GPU is more similar to a whole motherboard where your PCIe connection is a really really fast network connection. There's lots of faults to this analogy but this closer than just comparing TFLOPs.
Nvidia's moat has always been "CUDA". Quotes because even that is a messier term than most think (Cutlass, CuBLAS, cuDNN, CuTe, etc). The new cards are just capable of things the older ones aren't. Mix between hardware and software.
I know this isn't a great answer but there is none. You'll probably get some responses and many of them will have parts of the story but it's hard to paint a real good picture in a comment. There's no answer that is both good and short.
There's a 2x performance hit from the weird restriction on fp32 accumulation, plus the fact that 5090 has "fake" Blackwell (no tcgen05) which limits the size and throughput of matrix multiplication through the tensor cores.
Isn't 5090 FE (roughly 2500 USD in my country) pretty good FLOP value? 32 GB VRAM (and flash attention pushes it even faster compared to apple/mps relatively cheap "vram")
One of the reasons they removed NVLink from consumer cards (they supported it before). There’s also an issue with power consumption (1xB200 vs 10x5090)
Today, training in "low precision" probably means computing FP8 x FP8 -> FP32. The FP32 accumulation is still important, but otherwise yes this works, especially if we're talking about MXFP8 as supported on Blackwell [0].
What's less proven is a recipe using MXFP4 x MXFP4 -> FP32 compute, e.g. [1], which needs more involved techniques to work. But if you get it to work stably, that pathway is running at full throughput on 5090.
Interesting. My assumption was one of the innovations of DeepSeek and the modern GPT models was performing low precision pretraining rather than just finetuning further. I didn't realize you still need accumulation at a higher precision anyway
Even the large cloud AI services are focusing on this too, because it drives down the average "cost per query", or whatever you want to call it. For inference, arguably more even than training, the smaller and more efficient they can get it, the better their bottom line.
Curious what issues you were having. The kernel should compile natively if you pass nvcc the correct arch flags, although it probably won't take advantage of any new hardware features.
Definitely going to save this for later and come back to it after I get some more CUDA experience under my belt. It feels so nice right now making nice beautiful to use pipeline code w/ npp and some CUDA kernels here and there, the code is much faster than what it's replacing, but then I look at this guy getting down into the weeds of memory bank contention, prefetching, loop invariance, etc. Makes me feel like I'm playing with LEGO, I'm a little jealous.
The tip that Nsight can run on Mac over SSH is great, too. I've been capturing and viewing data over RDP, basically, will have to give it a shot next week.
> Due to improvements in newer hardware, you might need to use more tricks to reach Speed-of-Light on older GPUs e.g. pipeline shared memory to register memory data movements.
On the contrary, older GPUs are a lot easier to hit rooflines on. Newer GPUs run so fast that they keep adding new tricks to remove bottlenecks. Not to discount the author's work here but a 5090 is pretty bad on the FLOPs/memory bandwidth ratio so it's comparatively easier to get throttled by tensor cores there; on datacenter hardware your tensor cores are so fast that you'll hit limits that were glossed over here.
For example, using Ampere "mma" instructions won't cut it, because they compute a really small MMA and force your input to live in registers. You'll need TMA to get data into shared memory and wmma to do a matrix multiply out of them. At those speeds you will run into issues with dispatching instructions and computing addresses (and doing out-of-bounds calculation) fast enough that you will need to offload it to specialized hardware to keep up with the tensor cores.
My issue with upgrading to the 5090 for workstation ML use is that it both has higher TDP than the 4090 and it can only be limited to 70% power (not 50% like the 4090).
Hmm, but supposing the accelerated NVIDIA specific inference data types were available for Triton, then you would just use that? Why not contribute to Triton, they accept PRs? Like so what if you do free product ecosystem development for NVIDIA and giant corporations by contributing to Triton?
I mean, you can look at the most recent commit and see that the infrastructure is being built out for this right now (of course OpenAI doesn't care about sm_120, though).
By all means, the guy could have written the triton fixes he needs and NOT sent it up stream. It would still make more sense to do that! He’s obviously an expert, and I was sincerely wondering, why bother with the C++ stuff if he already knew the better way, and also has the chops to implement it?
I was surprised to see 5090's theoretical BF16 TFLOPs at just 209.5. That's not even 10% of the server Blackwell (B200 is 2250, and GB200 is 2500). B200 costs around $30-40k per GPU, so they are pretty close in performance per dollar.
Starting with 4090, NVIDIA limits the performance of tensor cores on gaming cards, specifically for ops that might be used in ML training. FP8 and FP16 matmuls run at full speed if accumulating in FP16 (I've never seen anyone use this), but only half speed when accumulating in FP32. This restriction is not present for lower precision matmuls like FP4, and is removed entirely on the workstation-class cards like RTX Pro 6000.
It doesn't seem worth it to use NVIDIA gaming cards as a "cheaper FLOPs" alternative anymore (e.g. diffusion models could have been cheaper to run on 3090 than A100). They are generous with memory bandwidth though, nearly 2TB/s on 5090 is amazing!
Is there really that big a different in TFLOPS between the GB100 and GB202 chips? The GB100 has fewer SMs than the GB202, so I'm confused about where the 10x performance would be coming from?
You're asking a really good question but it's not a question with an easy answer.
There's a lot more to performance computing than FLOPs. FLOPs are you good high level easy to understand metric but it's a small part of the story when you're in the weeds.
To help make sense of this, look at CPU frequencies. I think most people on HN know that two CPU with the same frequency can have dramatically different outcomes on benchmarks, right? You might know how some of these come down to things like IPC (instructions per cycle) or the cache structures. There's even more but we know it's not so easy to measure, right?
On a GPU all that is true but there's only more complexity. Your GPU is more similar to a whole motherboard where your PCIe connection is a really really fast network connection. There's lots of faults to this analogy but this closer than just comparing TFLOPs.
Nvidia's moat has always been "CUDA". Quotes because even that is a messier term than most think (Cutlass, CuBLAS, cuDNN, CuTe, etc). The new cards are just capable of things the older ones aren't. Mix between hardware and software.
I know this isn't a great answer but there is none. You'll probably get some responses and many of them will have parts of the story but it's hard to paint a real good picture in a comment. There's no answer that is both good and short.
No, GPUs are a lot simpler. You can mostly just take the clock rate and scale it directly for the instruction being compared.
There's a 2x performance hit from the weird restriction on fp32 accumulation, plus the fact that 5090 has "fake" Blackwell (no tcgen05) which limits the size and throughput of matrix multiplication through the tensor cores.
Isn't 5090 FE (roughly 2500 USD in my country) pretty good FLOP value? 32 GB VRAM (and flash attention pushes it even faster compared to apple/mps relatively cheap "vram")
Not really:
5090: 210 TF / $2k == 105 TF/$k
B200: 2250 TF / $40k == 56 TF/$k
Getting only 2x the FLOPs per dollar probably isn't worth the hassle of having to rack 10x as many GPUs, while having no NVLink.
One of the reasons they removed NVLink from consumer cards (they supported it before). There’s also an issue with power consumption (1xB200 vs 10x5090)
Sure, but when spending 20x more, getting almost twice the compute per buck seems expected
Isn't the new trend to train in lower precision anyway?
Today, training in "low precision" probably means computing FP8 x FP8 -> FP32. The FP32 accumulation is still important, but otherwise yes this works, especially if we're talking about MXFP8 as supported on Blackwell [0].
What's less proven is a recipe using MXFP4 x MXFP4 -> FP32 compute, e.g. [1], which needs more involved techniques to work. But if you get it to work stably, that pathway is running at full throughput on 5090.
[0]: https://arxiv.org/abs/2506.08027 [1]: https://arxiv.org/abs/2502.20586
Interesting. My assumption was one of the innovations of DeepSeek and the modern GPT models was performing low precision pretraining rather than just finetuning further. I didn't realize you still need accumulation at a higher precision anyway
Only GPU-poors run Q-GaLore and similar tricks.
Even the large cloud AI services are focusing on this too, because it drives down the average "cost per query", or whatever you want to call it. For inference, arguably more even than training, the smaller and more efficient they can get it, the better their bottom line.
For inference of course; the OP I replied to mentioned training though.
Do you have a source for that B200 price?
I had a 5090 some months ago but couldnt get flash attention to work. Does it now work natively? What about 5080?
Pytorch now has native support for the Blackwell architecture:
https://pytorch.org/blog/pytorch-2-7/
It does, but the performance is pretty bad, worse than Hopper.
Curious what issues you were having. The kernel should compile natively if you pass nvcc the correct arch flags, although it probably won't take advantage of any new hardware features.
High-performance GPU code typically uses nonportable features that are not supported across generations.
Definitely going to save this for later and come back to it after I get some more CUDA experience under my belt. It feels so nice right now making nice beautiful to use pipeline code w/ npp and some CUDA kernels here and there, the code is much faster than what it's replacing, but then I look at this guy getting down into the weeds of memory bank contention, prefetching, loop invariance, etc. Makes me feel like I'm playing with LEGO, I'm a little jealous.
The tip that Nsight can run on Mac over SSH is great, too. I've been capturing and viewing data over RDP, basically, will have to give it a shot next week.
Damn awesome. This going to take me 3 reads and a week to digest
> Due to improvements in newer hardware, you might need to use more tricks to reach Speed-of-Light on older GPUs e.g. pipeline shared memory to register memory data movements.
On the contrary, older GPUs are a lot easier to hit rooflines on. Newer GPUs run so fast that they keep adding new tricks to remove bottlenecks. Not to discount the author's work here but a 5090 is pretty bad on the FLOPs/memory bandwidth ratio so it's comparatively easier to get throttled by tensor cores there; on datacenter hardware your tensor cores are so fast that you'll hit limits that were glossed over here.
For example, using Ampere "mma" instructions won't cut it, because they compute a really small MMA and force your input to live in registers. You'll need TMA to get data into shared memory and wmma to do a matrix multiply out of them. At those speeds you will run into issues with dispatching instructions and computing addresses (and doing out-of-bounds calculation) fast enough that you will need to offload it to specialized hardware to keep up with the tensor cores.
My issue with upgrading to the 5090 for workstation ML use is that it both has higher TDP than the 4090 and it can only be limited to 70% power (not 50% like the 4090).
Hmm, but supposing the accelerated NVIDIA specific inference data types were available for Triton, then you would just use that? Why not contribute to Triton, they accept PRs? Like so what if you do free product ecosystem development for NVIDIA and giant corporations by contributing to Triton?
Second line of the post:
> 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.
Yes… I read it. If the feature is missing, why not contribute it instead?
How many PRs do you have landed in Triton that you can just blithely say "contribute it"?
I mean, you can look at the most recent commit and see that the infrastructure is being built out for this right now (of course OpenAI doesn't care about sm_120, though).
i don't know what this comment has to do with my point that OAI doesn't take commits from randoms, especially for infra code.
By all means, the guy could have written the triton fixes he needs and NOT sent it up stream. It would still make more sense to do that! He’s obviously an expert, and I was sincerely wondering, why bother with the C++ stuff if he already knew the better way, and also has the chops to implement it?
There's an enormous difference between writing kernels and writing compiler infra.
Yeah they do