[Openmp-commits] [PATCH] D64648: [OPENMP][NVPTX]Fixed checks for cuda versions.

Alexey Bataev via Phabricator via Openmp-commits openmp-commits at lists.llvm.org
Fri Jul 12 10:48:31 PDT 2019

ABataev created this revision.
ABataev added a reviewer: grokos.
Herald added subscribers: jdoerfert, guansong.
Herald added a project: OpenMP.

We used CUDART_VERSION macro to check for the installed cuda version
but this macro is defined in cuda_runtime_api.h, which is not used by
project. Better to use CUDA_VERSION macro, which is defined in cuda.h.
Also, added the check if this macro is defined. If macro is undefined,
there is something wrong with the cuda configuration and we should not
continue the compilation.
This also fixes problems with runtime building in cuda 10+.

  rOMP OpenMP



Index: libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
--- libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
+++ libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h
@@ -48,7 +48,9 @@
 // Macros for Cuda intrinsics
 // In Cuda 9.0, the *_sync() version takes an extra argument 'mask'.
 // Also, __ballot(1) in Cuda 8.0 is replaced with __activemask().
-#if defined(CUDART_VERSION) && CUDART_VERSION >= 9000
+#error CUDA_VERSION macro is undefined, something wrong with cuda.
+#elif CUDA_VERSION >= 9000
 #define __SHFL_SYNC(mask, var, srcLane) __shfl_sync((mask), (var), (srcLane))
 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down_sync((mask), (var), (delta), (width))
@@ -58,7 +60,7 @@
 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down((var), (delta), (width))
 #define __ACTIVEMASK() __ballot(1)
+#endif // CUDA_VERSION
 #define __SYNCTHREADS_N(n) asm volatile("bar.sync %0;" : : "r"(n) : "memory");
 // Use original __syncthreads if compiled by nvcc or clang >= 9.0.

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D64648.209528.patch
Type: text/x-patch
Size: 1170 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190712/b16160b0/attachment.bin>

More information about the Openmp-commits mailing list