Live data from Hacker News

Show HN: Flash Attention in ~100 lines of CUDA

github.com

11–20 of 44 posts

Re: Show HN: Flash Attention in ~100 lines of CUDA

#11
post #8

What 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.

Re: Show HN: Flash Attention in ~100 lines of CUDA

#12
post #3

For those who have no idea what's being discussed, quick background. Discussing: Transformer [1] memory issues and approximate attention [2] in machine learning training. Specifically: FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. [3] As a side comment, this entire industry is sorely in need of at least intros. The entire space has moved so fast in the last year I need an entire new dic…

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…

I personally have been looking for "explain it like I'm a CS PhD with lots of experience and the ability to look stuff up". But I suspect your summary would be pretty handy as well.

Re: Show HN: Flash Attention in ~100 lines of CUDA

#14

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.

> 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

#15
post #4

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.

You mean triton the inference server or triton the DSL for cuda?

triton the DSL.

Re: Show HN: Flash Attention in ~100 lines of CUDA

#16

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.

> 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.

Re: Show HN: Flash Attention in ~100 lines of CUDA

#17
post #3

For those who have no idea what's being discussed, quick background. Discussing: Transformer [1] memory issues and approximate attention [2] in machine learning training. Specifically: FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. [3] As a side comment, this entire industry is sorely in need of at least intros. The entire space has moved so fast in the last year I need an entire new dic…

Zero shot is wrong, but that definition is commonly used.

Zero shot is testing out if distribution, not just "a task" not trained on. The later is ill defined.

The original definition comes from a few papers. But the classic example is a clarifier recognizing zebras but having never been trained in zebras (but may have been trained on horses). There's are out of distribution. But importantly, out of the implicit distribution, not the target distribution.

The common improper usage usually confuses these two. A simple example might me training in 256x256 images and testing on 1024x1024. That's still in the implicit distribution (as long as the classes are identical). A very common example is training on a large dataset like LAION and then testing on coco or image net 1k. This is not zero shot because the classes in ImageNet are in LAION (and in Coco). Basically, this is a useless definition because then any validation or test set is zero shot because those were never seen in the training data and thus out of the training distribution. But remember that data sets are proxies for larger distributions.

Where is can get sometimes tricky is tasks (emergence has entered the chat). For example, you may not intend to train a generative model to do clarification but you probably did (it's very clear -- in the math -- if you're training density models (KLD, score, etc)). This can get hairy because it's very easy to train a model to do things that you aren't realizing you are and later find out. Some people can get upset about this but it's the nature of frameworks that have low interpretability. There's still a lot of mathematics we need to learn and it tends not to be an explicit focus in ML but there are plenty in the community focused on this.

Re: Show HN: Flash Attention in ~100 lines of CUDA

#18
post #3

For those who have no idea what's being discussed, quick background. Discussing: Transformer [1] memory issues and approximate attention [2] in machine learning training. Specifically: FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. [3] As a side comment, this entire industry is sorely in need of at least intros. The entire space has moved so fast in the last year I need an entire new dic…

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…

frankly i don’t really feel like all that much has changed since 2020 except for scale

Re: Show HN: Flash Attention in ~100 lines of CUDA

#20
post #3

For those who have no idea what's being discussed, quick background. Discussing: Transformer [1] memory issues and approximate attention [2] in machine learning training. Specifically: FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. [3] As a side comment, this entire industry is sorely in need of at least intros. The entire space has moved so fast in the last year I need an entire new dic…

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.

Post reply on HN