Re:Show HN: Flash Attention in ~100 lines of CUDA
by treesciencebot on Saturday March 16, 2024 @12:18PM
(#39727161)
Pretty neat implementation. In general, for these sort of exercises (and even if the intention is to go to prod with custom kernels) I lean towards Triton to write the kernels themselves. It is much more easier to integrate to the tool chain, and allows a level of abstraction that doesn't affect performance even a little bit while providing useful constructs.
[ Parent
| on HN ]
Re:Show HN: Flash Attention in ~100 lines of CUDA
by whimsicalism on Saturday March 16, 2024 @01:21PM
(#39727718)
yeah even the official flashattention is moving many implementations from cutlass to triton except for the main mha backward/forward pass
Re:Show HN: Flash Attention in ~100 lines of CUDA
by jart on Saturday March 16, 2024 @06:26PM
(#39730033)
It was written with cutlass? No wonder Peter Kim found it valuable and worthwhile to de-obfuscate. Adopting a new programming language invented by OpenAI doesn't sound like a much better alternative. I'd be shocked if either of them were able to build code for AMD GPUs, where it's easy to adapt CUDA code, but not if it's buried in tens of thousands of lines of frameworks. I like open source code to have clarity so I can optimize it for my own production environment myself. When people distribute code they've productionized for themselves, it squeezes out all the alpha and informational value. Just because something's open source doesn't mean it's open source. I think people mostly do it to lick the cookie without giving much away.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by synquid on Sunday March 17, 2024 @05:11AM
(#39733002)
Triton has an AMD backend, although work is still ongoing.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by imtringued on Sunday March 17, 2024 @08:26AM
(#39733986)
You will also be able to use Triton to target Ryzen AI.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by ixaxaar on Saturday March 16, 2024 @01:10PM
(#39727626)
You mean triton the inference server or triton the DSL for cuda?
Re:Show HN: Flash Attention in ~100 lines of CUDA
by p1esk on Saturday March 16, 2024 @01:21PM
(#39727715)
The DSL: https://openai.com/research/triton
Re:Show HN: Flash Attention in ~100 lines of CUDA
by whimsicalism on Saturday March 16, 2024 @01:21PM
(#39727713)
they mean the dsl (not just necessarily for cuda)
Re:Show HN: Flash Attention in ~100 lines of CUDA
by treesciencebot on Saturday March 16, 2024 @02:16PM
(#39728148)
triton the DSL.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by fpgamlirfanboy on Saturday March 16, 2024 @02:02PM
(#39728029)
> allows a level of abstraction that doesn't affect performance even a little bit
The second part of this sentence is true because the first part is false.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by treesciencebot on Saturday March 16, 2024 @02:17PM
(#39728155)
zero cost abstractions exist. doesn't mean all abstractions are zero-cost. or being zero-cost somehow invalidates their abstractness/genericness. but maybe we differ on the definition of abstractions.
Re:Show HN: Flash Attention in ~100 lines of CUDA
by queuebert on Saturday March 16, 2024 @08:35PM
(#39730758)
As a person who finds CUDA extremely easy to write and integrate, what does Triton have to offer?
Re:Show HN: Flash Attention in ~100 lines of CUDA
by whimsicalism on Saturday March 16, 2024 @10:09PM
(#39731325)
block level rather than thread level programming, automatic optimization across hyperparameters, makes it much easier to write fast kernels