CUDA __syncthreads() malfunctioning with -O2 optimization or higher (clang 14.0.6 and 15.0.3)

[cbuchner1]

I use CUDA 11.3 and an nVidia A100, hence the architecture sm_80 is specified during compilation.

__syncthreads() is no longer working for me when compiling CUDA code ever since upgrading from Clang 12 to Clang 14.0.6. The problem persists in Clang 15.0.3

I have a short reproducer. It throws runtime asserts when compiled with -O2 and above.

// clang++ -O3 --cuda-gpu-arch=sm_80 -x cuda test.cu -o test -L/usr/local/cuda-11.3/lib64 -lcudart
#include <cassert>
__global__ void test()
{
   __shared__ int test;
   test = 0;
   __syncthreads();
   if (threadIdx.x == 0 && threadIdx.y == 0 && threadIdx.z == 0)
   {
     test = 1234;
   }
   __syncthreads();
   assert(test == 1234);
}
#include <iostream>
int main(int argc, char **argv)
{
  dim3 block(16,16,1);
  dim3 grid(1,1,1);
  test<<<grid, block>>>();
  cudaDeviceSynchronize();
  std::cerr << "CUDA error code: " << cudaGetLastError() << std::endl;
}

As a workaround, when I call the below inline assembly function barSync() in place of __syncthreads(), the code starts working again.

// inline assembly to insert a barrier synchronization equivalent to __syncthreads()
__device__ __forceinline__ void barSync() {
      asm volatile("bar.sync 0;" : : : "memory");
}

[Artem-B]

This may be the same issue as these:
https://lists.llvm.org/pipermail/llvm-dev/2021-November/154060.html
#54851

It should be already fixed in HEAD.

[Artem-B]

Fixed by 9dc7da3 #54851