Show HN: Flash Attention in ~100 lines of CUDA

Re:Show HN: Flash Attention in ~100 lines of CUDA by saiojd on Saturday March 16, 2024 @01:28PM (#39727770)

What does __syncthreads() do here exactly? I'm new to CUDA, could get the overall idea of the FlashAttention paper but not the details.

[ Parent | on HN ]

Re:Show HN: Flash Attention in ~100 lines of CUDA by cavisne on Saturday March 16, 2024 @01:33PM (#39727804)

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.

[ Parent | on HN ]

Re:Show HN: Flash Attention in ~100 lines of CUDA by xrd on Saturday March 16, 2024 @04:43PM (#39729333)

Any CUDA primer you recommend in particular? I had this same question.

[ Parent | on HN ]

Re:Show HN: Flash Attention in ~100 lines of CUDA by winwang on Saturday March 16, 2024 @05:56PM (#39729842)

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

[ Parent | on HN ]

Re:Show HN: Flash Attention in ~100 lines of CUDA by cavisne on Saturday March 16, 2024 @07:02PM (#39730231)

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

[ Parent | on HN ]

Re:Show HN: Flash Attention in ~100 lines of CUDA by xrd on Saturday March 16, 2024 @07:38PM (#39730435)

This is terrific, thanks!

[ Parent | on HN ]