[libomptarget] Refactor syncwarp macro to inline function

[JonChesterfield]

Summary

[libomptarget] Refactor syncwarp macro to inline function
See also abandoned D66846, split into this diff and others.

[ABataev]

LG with a nit.

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

50

#endif // CUDA_VERSION

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 // Also, __ballot(1) in Cuda 8.0 is replaced with __activemask().
 #ifndef CUDA_VERSION
 #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))
 #define __ACTIVEMASK() __activemask()
-#define __SYNCWARP(Mask) __syncwarp(Mask)
 #else
 #define __SHFL_SYNC(mask, var, srcLane) __shfl((var), (srcLane))
 #define __SHFL_DOWN_SYNC(mask, var, delta, width)                              \
   __shfl_down((var), (delta), (width))
 #define __ACTIVEMASK() __ballot(1)
-// In Cuda < 9.0 no need to sync threads in warps.
-#define __SYNCWARP(Mask)
 #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.
 #if !defined(__clang__) || __clang_major__ >= 9
 #define __SYNCTHREADS() __syncthreads()
 #else
 #define __SYNCTHREADS() __SYNCTHREADS_N(0)

openmp/libomptarget/deviceRTLs/nvptx/src/supporti.h

 //===--------- supporti.h - NVPTX OpenMP support functions ------- CUDA -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper implementation to some functions natively supported by the GPU.
 //
 //===----------------------------------------------------------------------===//
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 
+#include "target_impl.h"
+
 INLINE void setExecutionParameters(ExecutionMode EMode, RuntimeMode RMode) {
   execution_param = EMode;
   execution_param |= RMode;
 }
 
 INLINE bool isGenericMode() { return (execution_param & ModeMask) == Generic; }
 
 INLINE bool isSPMDMode() { return (execution_param & ModeMask) == Spmd; }

 
 INLINE int IsTeamMaster(int ompThreadId) { return (ompThreadId == 0); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Parallel level
 
 INLINE void IncParallelLevel(bool ActiveParallel) {
   unsigned Active = __ACTIVEMASK();
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
   unsigned LaneMaskLt;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
   unsigned Rank = __popc(Active & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] +=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
 }
 
 INLINE void DecParallelLevel(bool ActiveParallel) {
   unsigned Active = __ACTIVEMASK();
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
   unsigned LaneMaskLt;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(LaneMaskLt));
   unsigned Rank = __popc(Active & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] -=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
     __threadfence();
   }
-  __SYNCWARP(Active);
+  __kmpc_impl_syncwarp(Active);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // get OpenMP number of procs
 
 // Get the number of processors in the device.
 INLINE int GetNumberOfProcsInDevice(bool isSPMDExecutionMode) {
   if (!isSPMDExecutionMode)

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h

INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(res));
   return res;
 }
 
 INLINE int __kmpc_impl_ffs(uint32_t x) { return __ffs(x); }
 
 INLINE int __kmpc_impl_popc(uint32_t x) { return __popc(x); }
 
-INLINE void __kmpc_impl_syncwarp(int32_t Mask) { __SYNCWARP(Mask); }
+#ifndef CUDA_VERSION
+#error CUDA_VERSION macro is undefined, something wrong with cuda.
+#endif
+
+INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
+#if CUDA_VERSION >= 9000
+  __syncwarp(Mask);
+#else
+  // In Cuda < 9.0 no need to sync threads in warps.
+#endif
+}
 
 #endif