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
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
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
