[clangd] Fix missing __syncthreads.

[oToToT]

Summary

I think the symbol for syncthreads is missing. Therefore, I've added this into the clang_cuda_device_functions.h.

This also solve an issue on GitHub opened about half a year ago.
Ref: https://github.com/clangd/clangd/issues/404

[oToToT]

I guess __nvvm_bar_sync is the right instruction.

[jdoerfert]

I doubt this compiles:

https://godbolt.org/z/c51v5d

even with adjusted type, this fails because __syncthreads is a builtin (right now):

https://godbolt.org/z/jP7e6h

(and __nvvm_bar_sync requires an argument).

[oToToT]

OK, I've found what's wrong with my fix here. I should check for the clangd compile mechanism with more detail to fix this problem.

Sorry for disturbance.

[jdoerfert]

In D96102#2545644, @oToToT wrote:

OK, I've found what's wrong with my fix here. I should check for the clangd compile mechanism with more detail to fix this problem.

Sorry for disturbance.

No worries. I'm not sure how clangd handles builtins, but I guess that is a general thing it should as this is not the only one people might want to use :)

clang/lib/Headers/__clang_cuda_device_functions.h

 }
 __DEVICE__ float __saturatef(float __a) { return __nv_saturatef(__a); }
 __DEVICE__ int __signbitd(double __a) { return __nv_signbitd(__a); }
 __DEVICE__ int __signbitf(float __a) { return __nv_signbitf(__a); }
 __DEVICE__ void __sincosf(float __a, float *__s, float *__c) {
   return __nv_fast_sincosf(__a, __s, __c);
 }
 __DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
+__DEVICE__ int __syncthreads() { return __nvvm_bar_sync(); }
 __DEVICE__ int __syncthreads_and(int __a) { return __nvvm_bar0_and(__a); }
 __DEVICE__ int __syncthreads_count(int __a) { return __nvvm_bar0_popc(__a); }
 __DEVICE__ int __syncthreads_or(int __a) { return __nvvm_bar0_or(__a); }
 __DEVICE__ float __tanf(float __a) { return __nv_fast_tanf(__a); }
 __DEVICE__ void __threadfence(void) { __nvvm_membar_gl(); }
 __DEVICE__ void __threadfence_block(void) { __nvvm_membar_cta(); };
 __DEVICE__ void __threadfence_system(void) { __nvvm_membar_sys(); };
 __DEVICE__ void __trap(void) { asm volatile("trap;"); }