[clang] f28e710 - [OpenMP] Make CUDA math library functions SPMD amenable

[Joseph Huber via cfe-commits]

Author: Joseph Huber
Date: 2021-09-10T14:52:45-04:00
New Revision: f28e710db720a913f4b508a9dc43f25e81629e72

URL: https://github.com/llvm/llvm-project/commit/f28e710db720a913f4b508a9dc43f25e81629e72
DIFF: https://github.com/llvm/llvm-project/commit/f28e710db720a913f4b508a9dc43f25e81629e72.diff

LOG: [OpenMP] Make CUDA math library functions SPMD amenable

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

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D108958

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_libdevice_declares.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_libdevice_declares.h b/clang/lib/Headers/__clang_cuda_libdevice_declares.h
index 6173b589e3eff..23f35964ea974 100644
--- a/clang/lib/Headers/__clang_cuda_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_cuda_libdevice_declares.h
@@ -16,6 +16,7 @@ extern "C" {
 
 #if defined(__OPENMP_NVPTX__)
 #define __DEVICE__
+#pragma omp begin assumes ext_spmd_amenable no_openmp
 #elif defined(__CUDA__)
 #define __DEVICE__ __device__
 #endif
@@ -456,6 +457,11 @@ __DEVICE__ double __nv_y1(double __a);
 __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 no_openmp
+#endif
+
 #if defined(__cplusplus)
 } // extern "C"
 #endif