[PATCH] D150894: [CUDA] provide wrapper functions for new NVCC builtins.

[Artem Belevich via Phabricator via cfe-commits]

tra updated this revision to Diff 523472.
tra added a comment.

Put the wrappers behind __CUDA_ARCH__ >= 800, as these clang builtins are not
available on older GPUs.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D150894

Files:
  clang/lib/Headers/__clang_cuda_intrinsics.h


Index: clang/lib/Headers/__clang_cuda_intrinsics.h
===================================================================
--- clang/lib/Headers/__clang_cuda_intrinsics.h
+++ clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -512,6 +512,68 @@
 __device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {
   return __nv_cvta_generic_to_shared_impl(__ptr);
 }
+
+#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800
+__device__ inline unsigned __reduce_add_sync_unsigned_impl(unsigned __mask,
+                                                           unsigned __value) {
+  return __nvvm_redux_sync_add(__mask, __value);
+}
+__device__ inline int __reduce_add_sync_signed_impl(unsigned __mask,
+                                                    int __value) {
+  return __nvvm_redux_sync_add(__mask, __value);
+}
+__device__ inline unsigned __reduce_min_sync_unsigned_impl(unsigned __mask,
+                                                           unsigned __value) {
+  return __nvvm_redux_sync_umin(__mask, __value);
+}
+__device__ inline unsigned __reduce_max_sync_unsigned_impl(unsigned __mask,
+                                                           unsigned __value) {
+  return __nvvm_redux_sync_umax(__mask, __value);
+}
+__device__ inline int __reduce_min_sync_signed_impl(unsigned __mask,
+                                                    int __value) {
+  return __nvvm_redux_sync_min(__mask, __value);
+}
+__device__ inline int __reduce_max_sync_signed_impl(unsigned __mask,
+                                                    int __value) {
+  return __nvvm_redux_sync_max(__mask, __value);
+}
+__device__ inline unsigned __reduce_or_sync_unsigned_impl(unsigned __mask,
+                                                          unsigned __value) {
+  return __nvvm_redux_sync_or(__mask, __value);
+}
+__device__ inline unsigned __reduce_and_sync_unsigned_impl(unsigned __mask,
+                                                           unsigned __value) {
+  return __nvvm_redux_sync_and(__mask, __value);
+}
+__device__ inline unsigned __reduce_xor_sync_unsigned_impl(unsigned __mask,
+                                                           unsigned __value) {
+  return __nvvm_redux_sync_xor(__mask, __value);
+}
+
+__device__ inline void
+__nv_memcpy_async_shared_global_4_impl(void *__dst, const void *__src,
+                                       unsigned __src_size) {
+  __nvvm_cp_async_ca_shared_global_4(
+      (void __attribute__((address_space(3))) *)__dst,
+      (const void __attribute__((address_space(1))) *)__src, __src_size);
+}
+__device__ inline void
+__nv_memcpy_async_shared_global_8_impl(void *__dst, const void *__src,
+                                       unsigned __src_size) {
+  __nvvm_cp_async_ca_shared_global_8(
+      (void __attribute__((address_space(3))) *)__dst,
+      (const void __attribute__((address_space(1))) *)__src, __src_size);
+}
+__device__ inline void
+__nv_memcpy_async_shared_global_16_impl(void *__dst, const void *__src,
+                                        unsigned __src_size) {
+  __nvvm_cp_async_ca_shared_global_16(
+      (void __attribute__((address_space(3))) *)__dst,
+      (const void __attribute__((address_space(1))) *)__src, __src_size);
+}
+#endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800
+
 } // extern "C"
 #endif // CUDA_VERSION >= 11000
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D150894.523472.patch
Type: text/x-patch
Size: 3360 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230518/de83985c/attachment-0001.bin>