Not Hacker News Logo

Not

Hacker

News!

Home
Hiring
Products
Companies
Discussion
Q&A
Users
Not Hacker News Logo

Not

Hacker

News!

AI-observed conversations & context

Daily AI-observed summaries, trends, and audience signals pulled from Hacker News so you can see the conversation before it hits your feed.

LiveBeta

Explore

  • Home
  • Hiring
  • Products
  • Companies
  • Discussion
  • Q&A

Resources

  • Visit Hacker News
  • HN API
  • Modal cronjobs
  • Meta Llama

Briefings

Inbox recaps on the loudest debates & under-the-radar launches.

Connect

© 2025 Not Hacker News! — independent Hacker News companion.

Not affiliated with Hacker News or Y Combinator. We simply enrich the public API with analytics.

Not Hacker News Logo

Not

Hacker

News!

Home
Hiring
Products
Companies
Discussion
Q&A
Users
  1. Home
  2. /Discussion
  3. /Writing Speed-of-Light Flash Attention for 5090 in CUDA C++
  1. Home
  2. /Discussion
  3. /Writing Speed-of-Light Flash Attention for 5090 in CUDA C++
Last activity 3 months agoPosted Aug 23, 2025 at 8:29 AM EDT

Writing Speed-of-Light Flash Attention for 5090 in Cuda C++

dsr12
159 points
34 comments

Mood

calm

Sentiment

positive

Category

other

Key topics

Cuda
Nvidia
GPU Optimization
Machine Learning
Debate intensity40/100

The post discusses optimizing flash attention for NVIDIA's 5090 GPU using CUDA C++, sparking a discussion on GPU performance, optimization techniques, and the trade-offs between different NVIDIA GPU models.

Snapshot generated from the HN discussion

Discussion Activity

Moderate engagement

First comment

3h

Peak period

8

Hour 15

Avg / period

2.6

Comment distribution34 data points
Loading chart...

Based on 34 loaded comments

Key moments

  1. 01Story posted

    Aug 23, 2025 at 8:29 AM EDT

    3 months ago

    Step 01
  2. 02First comment

    Aug 23, 2025 at 11:37 AM EDT

    3h after posting

    Step 02
  3. 03Peak activity

    8 comments in Hour 15

    Hottest window of the conversation

    Step 03
  4. 04Latest activity

    Aug 24, 2025 at 5:17 AM EDT

    3 months ago

    Step 04

Generating AI Summary...

Analyzing up to 500 comments to identify key contributors and discussion patterns

Discussion (34 comments)
Showing 34 comments
doctorpangloss
3 months ago
1 reply
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?
qeternity
3 months ago
1 reply
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.

doctorpangloss
3 months ago
1 reply
Yes… I read it. If the feature is missing, why not contribute it instead?
almostgotcaught
3 months ago
1 reply
How many PRs do you have landed in Triton that you can just blithely say "contribute it"?
saagarjha
3 months ago
1 reply
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).
almostgotcaught
3 months ago
2 replies
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.
doctorpangloss
3 months ago
1 reply
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?
almostgotcaught
3 months ago
There's an enormous difference between writing kernels and writing compiler infra.
saagarjha
3 months ago
Yeah they do
steinvakt2
3 months ago
2 replies
I had a 5090 some months ago but couldnt get flash attention to work. Does it now work natively? What about 5080?
sigmoid10
3 months ago
1 reply
Pytorch now has native support for the Blackwell architecture:

https://pytorch.org/blog/pytorch-2-7/

SynasterBeiter
3 months ago
It does, but the performance is pretty bad, worse than Hopper.
zackangelo
3 months ago
1 reply
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.
saagarjha
3 months ago
High-performance GPU code typically uses nonportable features that are not supported across generations.
neilmovva
3 months ago
4 replies
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!

steinvakt2
3 months ago
1 reply
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")
neilmovva
3 months ago
2 replies
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.

steinvakt2
3 months ago
Sure, but when spending 20x more, getting almost twice the compute per buck seems expected
lossolo
3 months ago
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)
mota7
3 months ago
2 replies
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?
godelski
3 months ago
1 reply
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.

saagarjha
3 months ago
No, GPUs are a lot simpler. You can mostly just take the clock rate and scale it directly for the instruction being compared.
saagarjha
3 months ago
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.
laidoffamazon
3 months ago
2 replies
Isn't the new trend to train in lower precision anyway?
storus
3 months ago
1 reply
Only GPU-poors run Q-GaLore and similar tricks.
Twirrim
3 months ago
1 reply
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.
storus
3 months ago
For inference of course; the OP I replied to mentioned training though.
neilmovva
3 months ago
1 reply
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

laidoffamazon
3 months ago
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
gautamcgoel
3 months ago
Do you have a source for that B200 price?
ProofHouse
3 months ago
Damn awesome. This going to take me 3 reads and a week to digest
saagarjha
3 months ago
> 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.

Scene_Cast2
3 months ago
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).
a_t48
3 months ago
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.

View full discussion on Hacker News
ID: 44995508Type: storyLast synced: 11/20/2025, 5:54:29 PM

Want the full context?

Jump to the original sources

Read the primary article or dive into the live Hacker News thread when you're ready.

Read ArticleView on HN
Not Hacker News Logo

Not

Hacker

News!

AI-observed conversations & context

Daily AI-observed summaries, trends, and audience signals pulled from Hacker News so you can see the conversation before it hits your feed.

LiveBeta

Explore

  • Home
  • Hiring
  • Products
  • Companies
  • Discussion
  • Q&A

Resources

  • Visit Hacker News
  • HN API
  • Modal cronjobs
  • Meta Llama

Briefings

Inbox recaps on the loudest debates & under-the-radar launches.

Connect

© 2025 Not Hacker News! — independent Hacker News companion.

Not affiliated with Hacker News or Y Combinator. We simply enrich the public API with analytics.