[PATCH] D96102: [clangd] Fix missing __syncthreads.

Tommy Chiang via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Thu Feb 4 21:08:15 PST 2021


oToToT updated this revision to Diff 321652.
oToToT added a comment.

I guess __nvvm_bar_sync is the right instruction.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D96102/new/

https://reviews.llvm.org/D96102

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() { 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); }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D96102.321652.patch
Type: text/x-patch
Size: 659 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210205/aadd5ee0/attachment.bin>


More information about the cfe-commits mailing list