CUDA : fix divergent FlashAttention barrier

[mistrjirka]

Overview

Incorrect use of __syncthreads() created divergent barriers in the FlashAttention code. It produced 1024 synchronization errors reported by NVIDIA Compute Sanitizer:
compute-sanitizer --tool synccheck.
The issue is:

.... threads in the same block  ... 
if (thread_type_a) {
    .....
    __syncthreads(); // creates barrier A
} else {
   __syncthreads();  // creates barrier B
}

So the threads in the same block are running towards different barriers based on which branch they take.
This can be fixed by:

.... threads in the same block ... 
if (thread_type_a) {
    .....
}

__syncthreads();  // creates common barrier

This produces 0 errors. The numerical tests passed and after benchmarking 100k cached + 1k prompt + 256 generated tokens produced almost no performance change.
Measured change: improved prompt processing by 0.43% and token generation by 0.15%, so no measurable performance regression.

Additional information

• https://docs.nvidia.com/cuda/cuda-programming-guide/#thread-block-synchronization-functions
• https://docs.nvidia.com/compute-sanitizer/ComputeSanitizer/index.html#synccheck-tool

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: yes, GPT-5.6 Sol. I was doing performance optimizations for qwen 3.8 27B and it identified this issue during some FlashAttention work.