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

Show HN: Flash Attention in ~100 lines of CUDA

github.com

Post reply on HN
Print view View original

31–40 of 44 posts

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

#31

Post by jhanoncomm »

jprete wrote: post #12
robrenaud wrote: post #10

Earlier 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…

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.

I reckon you need tacit knowledge. Experience. Luckily in the order of 100 hours not 10000.

Build a GPT using Python and Pytorch. For a good course: Andrej Karpathy is your keyword. At $1000 his course is great value. But actually it is free which is even better ;-)

It wont take you to flash attention but will ramp you to the point you could probably read papers about it. I almost got that far then life lifed me. But I was able to implement changes to the architecture of GPT and do some “hey mum I am doing SOTA (2021) machine learning”.

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

#32

Post by cavisne »

xrd wrote: post #26
cavisne wrote: post #11

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

Probably https://www.youtube.com/watch?v=nOxKexn3iBo (or just skimming the attached colab).

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

#34

Post by queuebert »

treesciencebot wrote: post #2

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.

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

#35

Post by andoando »

robrenaud wrote: post #10
araes wrote: 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 would love an explanation for software enginners / CS majors who aren't familiar with ML.

Last I studied ML was 2016 and that was stuff like decision trees, k nearest neighbors...

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

#36

Post by chillee »

einpoklum wrote: post #29

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

> maybe this is useful and maybe it isn't, but counting lines of code on top of a huge codebase is not very meaningful.

In this case it's pretty reasonable imo, since the kernel itself is fairly independent - the usage of torch is just for some bindings for the data structures.

> Launching separate kernels, synchronously, on the default stream, for various operations, is typically not the right way to utilize a GPU.

This is actually the standard way to do things in ML. Assuming you're from a HPC background (where this may seem quite strange), the biggest change is that "More or less everything in ML runs on the GPU", so there is very rarely any device to host synchronizations. In addition, each individual kernel is typically run on fairly large chunks of data (a million elements would be on the smaller side), so maximizing occupancy with streams is not as necessary as in HPC.

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

#37

Post by whimsicalism »

queuebert wrote: post #34
treesciencebot wrote: post #2

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.

As a person who finds CUDA extremely easy to write and integrate, what does Triton have to offer?

block level rather than thread level programming, automatic optimization across hyperparameters, makes it much easier to write fast kernels

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

#40

Post by synquid »

jart wrote: post #30
whimsicalism wrote: post #7

Earlier quoted context omitted.

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…

Triton has an AMD backend, although work is still ongoing.
Post reply on HN