Hacker News Board Hacker News, browsed like a classic message board
Live data from Hacker News

FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

together.ai

Post reply on HN
Print view View original

31–40 of 63 posts

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#31

Post by kristjansson »

refulgentis wrote: post #23
kristjansson wrote: post #3

Earlier quoted context omitted.

FlashAttention's algorithmic improvements is mostly just splitting/combining the softmax part of attention, and is itself not totally novel. The overwhelming contribution is implementing that, and all its fiddly pieces, efficiently on Nvidia hardware.

Clarifying: Given the question: "How much is the flash attention algorithm tied to the hardware?" The answer is 0. ex. you can find generic flash attention recently added in llama.cpp and ONNX (MS needed it for Phi-3, needed for Recall). On the side, novelty, I have no direct knowledge on, IMHO, asking that question would devolve the way novelty arguments do in any field: there's always someone else who can claim the…

I think the right analogy for FA is high-quality cache-aware BLAS kernel implementations. The algorithm(s) is (are) clever and (as you note) completely independent of hardware. However, a hardware-naive implementation is approximately worthless. Most of the value of MKL, or Accelerate, or FA is in the careful matching of the parameters and implementation of the algorithm to the capabilities of hardware it's going run on.

I definitely don't mean to take away from Tri/FA by mentioning novelty - I'm just repeating from paper, which refers back to algebraic aggregates[0] in its discussion of their tiled softmax.

[0]: https://web.stanford.edu/class/cs345d-01/rl/olap.pdf

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#32

Post by kristjansson »

f_devd wrote: post #5
edude03 wrote: post #2

How much is the flash attention algorithm tied to the hardware? For example, in this announcement they mention taking advantage of the async capabilities of the H100 GPUs which I assume means you won't get those speedups on non H series card. Two, the actual flash attention library requires CUDA, although the algorithm has apparently?[^0] been ported to metal. I would imagine if the algorithm was literally just a pur…

> How much is the flash attention algorithm tied to the hardware? The original FA, almost none. For the latest versions depends on your abstraction, ThunderKittens[0] provides about the same speed up over FA2 (1.3x-2x%) as the article but relatively universal across GPUs. For any new hardware there may be hardware specific features that make it edge out more performance; usually vendors will adopt any new features th…

I mean they're building an API to abstract away some of the SKU-to-SKU differences, but the broader point cuts the other way, I think:

> In fact, more broadly we believe we should really reorient our ideas of AI around what maps well onto the hardware. How big should a recurrent state be? As big can fit onto an SM. How dense should the compute be? No less so than what the hardware demands. An important future direction of this work for us is to use our learnings about the hardware to help us design the AI to match.

The value is in adapting the implementation (either manually at write-time or programmatically at run-time) to the specifics of the hardware.

Also, great line:

> And we ask: if your matrix multiply is smaller than 16x16, are you sure what you’re doing is AI?

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#34

Post by Lerc »

WanderPanda wrote: post #14

Compiler folks: Is there any chance compilers will be able to find optimizations like FlashAttention on their own? Seems like TVM and tinygrad are working in that direction but I find it hard to believe that that would be feasible

This strikes me as an extremely difficult but not intractable problem.

I'm not sure what the state of the art in compiler optimisation is with regard to data positioning and targeting maximum processor usage

There was a video on optimisation a while back that showed small optimisations caused increases in speed that were insignificant when compared to the speed variance induced by the memory layout that the optimisation (or even a random change) caused.

While that talk was more focused on getting a signal past the noise. That noise itself is an artifact of compilers being not particularly good at handling a much simpler form of the problem you describe.

CPU and memory architectures are complex when caches and access patterns impact upon speed.

When you add in GPU architectures to the mix I think you might be in fairly uncharted territory.

Maybe one day.

Of course since we are in the field of AI there is also the question of could a sufficiently smart AI do this. It depends on the value of sufficient.

I would like to think that an extremely high level test for an AI model could be to give it something like micrograd and tell it to produce something with the same interface that outperforms torch.

We're not even in the ballpark of being able to do that yet, but it will be interesting when and if that happens.

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#35

Post by eyegor »

f_devd wrote: post #5
edude03 wrote: post #2

How much is the flash attention algorithm tied to the hardware? For example, in this announcement they mention taking advantage of the async capabilities of the H100 GPUs which I assume means you won't get those speedups on non H series card. Two, the actual flash attention library requires CUDA, although the algorithm has apparently?[^0] been ported to metal. I would imagine if the algorithm was literally just a pur…

> How much is the flash attention algorithm tied to the hardware? The original FA, almost none. For the latest versions depends on your abstraction, ThunderKittens[0] provides about the same speed up over FA2 (1.3x-2x%) as the article but relatively universal across GPUs. For any new hardware there may be hardware specific features that make it edge out more performance; usually vendors will adopt any new features th…

What do you mean by "relatively universal"? This is Cuda only [0] with a promise of a rocm backend eventually. There's only one project I'm aware of that seriously tries to address the Cuda issue in ml [1].

[0] https://github.com/HazyResearch/ThunderKittens?tab=readme-ov...

[1] https://github.com/vosen/ZLUDA

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#36

Post by kiratp »

kristjansson wrote: post #31
refulgentis wrote: post #23

Earlier quoted context omitted.

Clarifying: Given the question: "How much is the flash attention algorithm tied to the hardware?" The answer is 0. ex. you can find generic flash attention recently added in llama.cpp and ONNX (MS needed it for Phi-3, needed for Recall). On the side, novelty, I have no direct knowledge on, IMHO, asking that question would devolve the way novelty arguments do in any field: there's always someone else who can claim the…

I think the right analogy for FA is high-quality cache-aware BLAS kernel implementations. The algorithm(s) is (are) clever and (as you note) completely independent of hardware. However, a hardware-naive implementation is approximately worthless. Most of the value of MKL, or Accelerate, or FA is in the careful matching of the parameters and implementation of the algorithm to the capabilities of hardware it's going run…

> However, a hardware-naive implementation is approximately worthless.

This isn’t true when there is one vendor that’s 90% of the market and 2 maybe 3 generations of hardware to consider. Support A100, H100 and you are supporting most of the current market.

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#38

Post by refibrillator »

The code has a comment which seems to hint that Tri Dao was working on FA3 as early as April 2022, the month after Hopper/H100 was announced. I find it mildly curious that over 2 years has elapsed before the code was released today. Perhaps it’s because now there’s better solutions in the pipeline?

Tri’s publication history has been leaning toward SSM and Mamba style architectures recently. Unlike Flash Attention which has quadratic time complexity wrt sequence length, these latest algorithms are subquadratic. Thus they do much less computation, instead of just doing it more efficiently a la Flash Attention.

Dao and Gu published a really long paper this year which demonstrated (among other things) how Mamba/SSM can be formulated such that it’s amenable to acceleration using the same hardware primitives that Transformers benefit from.

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#39

Post by JackYoustra »

latchkey wrote: post #37

If anyone wants to port this over to ROCm / AMD MI300x, reach out to me: hello@hotaisle.xyz (we won't ever spam you). Happy to donate the compute time for this work.

You're the AMD accelerator server company! Such cool work, hope someone takes you up :)

Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision

#40

Post by slashdave »

WanderPanda wrote: post #14

Compiler folks: Is there any chance compilers will be able to find optimizations like FlashAttention on their own? Seems like TVM and tinygrad are working in that direction but I find it hard to believe that that would be feasible

Kinda tricky if you want to call higher level operators in a wrapped language like Python.
Post reply on HN