Do no use bit_cast to work around initialization issues with barrier

None of the tests is constexpr so we can just use `reinterpret_cast`

Fixes #6255

Author: miscco

libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk.pass.cpp

@@ -43,7 +43,7 @@ __device__ void test()
   alignas(16) __shared__ int smem_buffer[buf_len];
 #if _CCCL_CUDA_COMPILER(CLANG)
   __shared__ char barrier_data[sizeof(barrier)];
-  barrier& bar = cuda::std::bit_cast<barrier>(barrier_data);
+  barrier& bar = reinterpret_cast<barrier&>(barrier_data);
 #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG)
   __shared__ barrier bar;
 #endif // !_CCCL_CUDA_COMPILER(CLANG)

libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_ptx_compiles.pass.cpp

@@ -30,7 +30,7 @@ __global__ void test_bulk_tensor(CUtensorMap* map)
   __shared__ int smem;
 #if _CCCL_CUDA_COMPILER(CLANG)
   __shared__ char barrier_data[sizeof(barrier)];
-  barrier& bar = cuda::std::bit_cast<barrier>(barrier_data);
+  barrier& bar = reinterpret_cast<barrier&>(barrier_data);
 #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG)
   __shared__ barrier bar;
 #endif // !_CCCL_CUDA_COMPILER(CLANG)
@@ -57,7 +57,7 @@ __global__ void test_bulk(void* gmem)
 {
   __shared__ int smem;
   __shared__ char barrier_data[sizeof(barrier)];
-  barrier& bar = *reinterpret_cast<barrier*>(&barrier_data);
+  barrier& bar = reinterpret_cast<barrier&>(barrier_data);
   if (threadIdx.x == 0)
   {
     init(&bar, blockDim.x);

libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor.pass.cpp

@@ -70,7 +70,7 @@ __device__ void test(int base_i, int base_j)
   alignas(128) __shared__ int smem_buffer[buf_len];
 #if _CCCL_CUDA_COMPILER(CLANG)
   __shared__ char barrier_data[sizeof(barrier)];
-  barrier& bar = cuda::std::bit_cast<barrier>(barrier_data);
+  barrier& bar = reinterpret_cast<barrier&>(barrier_data);
 #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG)
   __shared__ barrier bar;
 #endif // !_CCCL_CUDA_COMPILER(CLANG)

libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor_generic.h

@@ -187,7 +187,7 @@ test(cuda::std::array<uint32_t, num_dims> smem_coord,
   alignas(128) __shared__ int smem_buffer[smem_len];
 #if _CCCL_CUDA_COMPILER(CLANG)
   __shared__ char barrier_data[sizeof(barrier)];
-  barrier& bar = cuda::std::bit_cast<barrier>(barrier_data);
+  barrier& bar = reinterpret_cast<barrier&>(barrier_data);
 #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG)
   __shared__ barrier bar;
 #endif // !_CCCL_CUDA_COMPILER(CLANG)