6 ms·

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

• cavisne 3y ago
  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.
  • xrd 3y ago
    Any CUDA primer you recommend in particular? I had this same question.
    • winwang 3y ago
      Here's an article on syncing in CUDA via cooperative groups: https://developer.nvidia.com/blog/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-primitives/ https://developer.nvidia.com/blog/using-cuda-warp-level-prim...
    • cavisne 3y ago
      Probably https://www.youtube.com/watch?v=nOxKexn3iBo https://www.youtube.com/watch?v=nOxKexn3iBo (or just skimming the attached colab).
      • xrd 3y ago
        This is terrific, thanks!