diff --git a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk.pass.cpp b/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk.pass.cpp index 6b4e640b10d..2129a4753b8 100644 --- a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk.pass.cpp +++ b/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_data); + barrier& bar = reinterpret_cast(barrier_data); #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG) __shared__ barrier bar; #endif // !_CCCL_CUDA_COMPILER(CLANG) diff --git a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_ptx_compiles.pass.cpp b/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_ptx_compiles.pass.cpp index e5b787cfe21..32153434e77 100644 --- a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_ptx_compiles.pass.cpp +++ b/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_data); + barrier& bar = reinterpret_cast(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_data); + barrier& bar = reinterpret_cast(barrier_data); if (threadIdx.x == 0) { init(&bar, blockDim.x); diff --git a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor.pass.cpp b/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor.pass.cpp index dba372bc639..d20fc1aac45 100644 --- a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor.pass.cpp +++ b/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_data); + barrier& bar = reinterpret_cast(barrier_data); #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG) __shared__ barrier bar; #endif // !_CCCL_CUDA_COMPILER(CLANG) diff --git a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor_generic.h b/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor_generic.h index 9232df6e311..12f40d05636 100644 --- a/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor_generic.h +++ b/libcudacxx/test/libcudacxx/cuda/barrier/cp_async_bulk_tensor_generic.h @@ -187,7 +187,7 @@ test(cuda::std::array 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_data); + barrier& bar = reinterpret_cast(barrier_data); #else // ^^^ _CCCL_CUDA_COMPILER(CLANG) ^^^ / vvv !_CCCL_CUDA_COMPILER(CLANG) __shared__ barrier bar; #endif // !_CCCL_CUDA_COMPILER(CLANG)