This is fantastic. I am just starting in the ML space (compile from compilers) and I love short kernels that I can use to understand things better with.
> compile from compilers
What does that mean?
Show HN: Flash Attention in ~100 lines of CUDA
21–30 of 44 posts
Re: Show HN: Flash Attention in ~100 lines of CUDA
#22Earlier quoted context omitted.
> 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.
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.
> zero cost abstractions exist
So does perpetual motion :shrug: but my point is Triton is not an abstraction in the least. Source: 1) I spent 6 months investigating targeting other backends 2) Phil himself said he doesn't care to support other backends https://github.com/openai/triton/pull/1797#issuecomment-1730...
Re: Show HN: Flash Attention in ~100 lines of CUDA
#23Earlier quoted context omitted.
Regarding your comment about how fast the research and industry is moving, would HN readers be interested in relevant one or two paragraph summaries that are basically "explain it like I am a machine learning engineer from 2020" but also knows the power of these models from a perspective of using ChatGPT or MS Copilot? That is, assume a fair amount of technical knowledge about the fundamentals, but don't assume that…
That sounds at least somewhat helpful. Honestly, a gradient for some of this stuff would be nice. Explain to me like I'm: "five", "a high schooler", "a college grad (not CS/ML/Eng)", "a CS/Eng not ML". Although in a couple years, kids in restaurants will probably telling me how they're leveling up attention on their neuro-pet. The singularity is steep.
singularity implies AI increases exponentially, not human intelligence. Kids will not talk about neural nets any time soon.
Re: Show HN: Flash Attention in ~100 lines of CUDA
#24Earlier quoted context omitted.
> 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.
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.
[deleted]
Re: Show HN: Flash Attention in ~100 lines of CUDA
#25Earlier quoted context omitted.
> 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.
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.
It's amazing how heavily provided hn is. I have a response here that's been deleted that is like 15 words, including a link to source that corroborates my claim but that response contains a transcribed emoji and so it's been deleted by dang or whomever. Lol super rich environment for discourse we've got going here.
Re: Show HN: Flash Attention in ~100 lines of CUDA
#26What does __syncthreads() do here exactly? I'm new to CUDA, could get the overall idea of the FlashAttention paper but not the details.
Causes every thread in the block to wait until they have reached this point. Worth reading a cuda primer for more details on blocks/warps. Since the threads are relying on each other to fill the SRAM with all needed data if you didn’t wait then values would be missing.
Any CUDA primer you recommend in particular? I had this same question.
Re: Show HN: Flash Attention in ~100 lines of CUDA
#27Earlier quoted context omitted.
Causes every thread in the block to wait until they have reached this point. Worth reading a cuda primer for more details on blocks/warps. Since the threads are relying on each other to fill the SRAM with all needed data if you didn’t wait then values would be missing.
Any CUDA primer you recommend in particular? I had this same question.
Here's an article on syncing in CUDA via cooperative groups: https://developer.nvidia.com/blog/cooperative-groups/
There's also explicit warp synchronization, i.e. __syncwarp(). More on warp primitives here: https://developer.nvidia.com/blog/using-cuda-warp-level-prim...
Re: Show HN: Flash Attention in ~100 lines of CUDA
#28Re: Show HN: Flash Attention in ~100 lines of CUDA
#29My GPU work is not in ML (deep or otherwise); but ...
1. "100 lines of CUDA" + PyTorch; maybe this is useful and maybe it isn't, but counting lines of code on top of a huge codebase is not very meaningful.
2. Launching separate kernels, synchronously, on the default stream, for various operations, is typically not the right way to utilize a GPU.
Re: Show HN: Flash Attention in ~100 lines of CUDA
#30Pretty 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.
yeah even the official flashattention is moving many implementations from cutlass to triton except for the main mha backward/forward pass
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.