I am wondering why flash attention is like 5x slower with variable masking than without it? Lack of good masking support almost zeros out the optimizations
Where are you seeing these benchmarks?
FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
41–50 of 63 posts
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#42> TMA (Tensor Memory Accelerator). This is a special hardware unit that accelerates the transfer of data between global memory and shared memory, taking care of all index calculation and out-of-bound predication. This frees up registers, which is a valuable resource to increase tile size and efficiency.
My understanding was that while it frees up registers it more importantly lets the hardware handle address generation, which can become a bottleneck as other operations around it become faster.
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#43Earlier quoted context omitted.
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.
Supporting A100 and H100 is the opposite of being hardware naive, though.
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#44How 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…
To add to the discussion, from a practical perspective, AMD hardware totally sucks and yet to have proper implementation with flash-attention-2. ROCm is moving to usable slowly, but not close to being even comparable with cuda.
Whi os it so hard to port FA2 to the m1300 instinct?
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#45If 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.
Not trying to be rude but what is the thinking behind this offer? Why would someone do this port for…free save for access to the hardware? What’s the upside for them?
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#46I was wondering... this post mentions that ops like sigmoid are very slow.
A lot of modern LLMs use activation functions with sigmoid or soft max like SiLU, Swish, and SOLU.
Does Relu take less of a performance hit, and if so, maybe it'd be better to go back to good old relu?
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#47If 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.
Not trying to be rude but what is the thinking behind this offer? Why would someone do this port for…free save for access to the hardware? What’s the upside for them?
Not a rude question. I'm building public HPC super computers, currently focused on AMD hardware. The one I'm about to deploy is Top 150, which is a pretty good start.
The goal is to encourage a developer flywheel. The more developers working with AMD hardware, the more hardware that is needed, the more hardware I can justify buying, the bigger my super computers get.
Nvidia has been doing the flywheel for years and it has clearly worked. Why not do the same for AMD? As I said in another thread, anyone who thinks that there should be a single provider for all of AI compute needs, will be on the wrong side of history.
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#48This is one of the most important improvements in all of AI, because it benefits most AI users by giving them access to more, faster, for the same hardware with little to no tradeoffs.
...for all those users with H100s.
Indeed.
Anyone who is doing anything important or at scale would be at least renting those, or even using an abstracted service that is on top of another service.
Those cost savings allow people to train things for cheaper, causing those cost savings to benefit almost everyone doing important stuff in the space.
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#49Earlier quoted context omitted.
To add to the discussion, from a practical perspective, AMD hardware totally sucks and yet to have proper implementation with flash-attention-2. ROCm is moving to usable slowly, but not close to being even comparable with cuda.
Whi os it so hard to port FA2 to the m1300 instinct?
Re: FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision
#50Earlier quoted context omitted.
> 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
If you read the article I linked they show that it's entirely based on 16x16 matrices (or "tiles") which is fairly standard across gpus.