[PATCH] D80464: [CUDA] Missing __syncthreads intrinsic in __clang_cuda_device_functions.h

[Boris Staletic via Phabricator via cfe-commits]

bstaletic created this revision.
bstaletic added a reviewer: aprantl.
Herald added subscribers: cfe-commits, yaxunl.
Herald added a project: clang.

Seems like the `__syncthreads` is missing from the 
`clang/lib/Headers/__clang_cuda_device_functions.h` file. To be honest, I don't 
know much about CUDA. This issue was noticed by a YouCompleteMe user who then 
made a pull request:

https://github.com/ycm-core/ycmd/pull/1438

I did not create any tests, because a similar patch did not include tests:

https://reviews.llvm.org/D43602


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D80464

Files:
  clang/lib/Headers/__clang_cuda_device_functions.h


Index: clang/lib/Headers/__clang_cuda_device_functions.h
===================================================================
--- clang/lib/Headers/__clang_cuda_device_functions.h
+++ clang/lib/Headers/__clang_cuda_device_functions.h
@@ -519,6 +519,7 @@
   return __nv_fast_sincosf(__a, __s, __c);
 }
 __DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
+__DEVICE__ int __syncthreads(void) { return __nvvm_bar0(); }
 __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); }



Index: clang/lib/Headers/__clang_cuda_device_functions.h
===================================================================
--- clang/lib/Headers/__clang_cuda_device_functions.h
+++ clang/lib/Headers/__clang_cuda_device_functions.h
@@ -519,6 +519,7 @@
   return __nv_fast_sincosf(__a, __s, __c);
 }
 __DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
+__DEVICE__ int __syncthreads(void) { return __nvvm_bar0(); }
 __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); }

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits