[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