[OpenMP] Make CUDA math library functions SPMD amenable

[jhuber6]

Summary

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

[jdoerfert]

LG

[jhuber6]

Adding no_openmp assumption as well.

[jhuber6]

Fixing end assumes using wrong syntax.

clang/lib/Headers/__clang_cuda_libdevice_declares.h

 #define __CLANG_CUDA_LIBDEVICE_DECLARES_H__
 
 #if defined(__cplusplus)
 extern "C" {
 #endif
 
 #if defined(__OPENMP_NVPTX__)
 #define __DEVICE__
+#pragma omp begin assumes ext_spmd_amenable no_openmp
 #elif defined(__CUDA__)
 #define __DEVICE__ __device__
 #endif
 
 __DEVICE__ int __nv_abs(int __a);
 __DEVICE__ double __nv_acos(double __a);
 __DEVICE__ float __nv_acosf(float __a);
 __DEVICE__ double __nv_acosh(double __a);

 __DEVICE__ int __nv_vsubus4(int __a, int __b);
 #endif  // CUDA_VERSION
 __DEVICE__ double __nv_y0(double __a);
 __DEVICE__ float __nv_y0f(float __a);
 __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, omp_no_openmp
+#endif
+
 #if defined(__cplusplus)
 } // extern "C"
 #endif
 #endif // __CLANG_CUDA_LIBDEVICE_DECLARES_H__