Cub can't be compiled on CLANG with CUDA 11.8

[ShuaiShao93]

After upgrading to CUDA11.8, the cub can't be compiled with error

/local/user/cuda/include/cub/block/../detail/device_synchronize.cuh:54:14: error: no matching function for call to '__cudaDeviceSynchronizeDeprecationAvoidance'
    result = __cudaDeviceSynchronizeDeprecationAvoidance();
    ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

/local/usr/cuda/include/cuda_device_runtime_api.h:141:60: note: candidate function not viable: call to __device__ function from __host__ function
    extern __device__ __cudart_builtin__ cudaError_t CUDARTAPI __cudaDeviceSynchronizeDeprecationAvoidance(void); ^ 1 error generated when compiling for sm_61.

By digging into the source code, it seems to be a bug from CUB and CUDA. In include/cub/block/../detail/device_synchronize.cuh, device_synchronize() is a host function decorated by CUB_EXEC_CHECK_DISABLE, but it seems with #411, CUB_EXEC_CHECK_DISABLE is no-op for clang?

[gevtushenko]

Hello @ShuaiShao93 and thank you for reporting this!
Just to make sure, are you using CUB from CUDA Toolkit 11.8?
Can you provide a reproducer?

[ShuaiShao93]

Hello @ShuaiShao93 and thank you for reporting this! Just to make sure, are you using CUB from CUDA Toolkit 11.8?

Yes.

Can you provide a reproducer?

We have our custom setup to build tensorflow with bazel + clang, so I'm not sure if there is a way to easily reproduce this. But I guess if you can sudo apt install cuda-11-8 and use clang to build //tensorflow/python or any library that depends on device_synchronize.cuh, or even a dummy c++ file that includes device_synchronize.cuh, it will fail.

[ShuaiShao93]

This is because we mistakenly define CUDACC for clang.