[PATCH] D108958: [OpenMP] Make CUDA math library functions SPMD amenable

[Joseph Huber via Phabricator via cfe-commits]

jhuber6 created this revision.
jhuber6 added a reviewer: jdoerfert.
Herald added subscribers: guansong, yaxunl.
jhuber6 requested review of this revision.
Herald added subscribers: cfe-commits, sstefan1.
Herald added a project: clang.

This patch adds the SPMD amenable assumption to the CUDA math library
defintions in Clang. Previously these functions would block SPMD
execution on the device because they're intrinsic calls into the library
and can't be calculated. These functions don't have side-effects so they
are safe to execute in SPMD mode.

Depends on D105937 <https://reviews.llvm.org/D105937>


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D108958

Files:
  clang/lib/Headers/__clang_cuda_libdevice_declares.h


Index: clang/lib/Headers/__clang_cuda_libdevice_declares.h
===================================================================
--- clang/lib/Headers/__clang_cuda_libdevice_declares.h
+++ clang/lib/Headers/__clang_cuda_libdevice_declares.h
@@ -16,6 +16,7 @@
 
 #if defined(__OPENMP_NVPTX__)
 #define __DEVICE__
+#pragma omp begin assumes ext_spmd_amenable
 #elif defined(__CUDA__)
 #define __DEVICE__ __device__
 #endif
@@ -456,6 +457,11 @@
 __DEVICE__ float __nv_y1f(float __a);
 __DEVICE__ float __nv_ynf(int __a, float __b);
 __DEVICE__ double __nv_yn(int __a, double __b);
+
+#if defined(__OPENMP_NVPTX__)
+#pragma omp end assumes ext_spmd_amenable
+#endif
+
 #if defined(__cplusplus)
 } // extern "C"
 #endif


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D108958.369583.patch
Type: text/x-patch
Size: 713 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20210830/96395727/attachment.bin>