[libomptarget] Refactor syncthreads macro to inline function

[JonChesterfield]

Summary

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

[ABataev]

openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu

556

Remove FIXME

571

And here too

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

129

Remove FIXME

172

Same here

openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu

78

Same

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

51

#endif // __clang__

[JonChesterfield]

Thanks. I wasn't sure whether to leave the FIXMEs in place or not, will drop them.

[JonChesterfield]

• remove fixme, add comment, rebase on master

[JonChesterfield]

Phabricator has decided to mix this patch in with the other patches that have landed on master since initially creating it. Presumably my mistake for rebasing. I'll abandon this and create a new diff. Sorry for the noise.

openmp/libomptarget/deviceRTLs/nvptx/src/data_sharing.cu

 //===----- data_sharing.cu - NVPTX OpenMP debug utilities -------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the implementation of data sharing environments/
 //
 //===----------------------------------------------------------------------===//
 #include "omptarget-nvptx.h"
+#include "target_impl.h"
+
 #include <stdio.h>
 
 // Warp ID in the CUDA block
 INLINE static unsigned getWarpId() { return threadIdx.x / WARPSIZE; }
 // Lane ID in the CUDA warp.
 INLINE static unsigned getLaneId() { return threadIdx.x % WARPSIZE; }
 
 // Return true if this is the first active thread in the warp.
EXTERN void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode,
   if (is_shared) {
     *frame = buf;
     return;
   }
   if (isSPMDExecutionMode) {
     if (GetThreadIdInBlock() == 0) {
       *frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
     }
     // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
-    __SYNCTHREADS();
+    __kmpc_impl_syncthreads();
     return;
   }
   ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   *frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
   __threadfence();
 }
 
 EXTERN void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
                                               int16_t is_shared) {
   if (is_shared)
     return;
   if (isSPMDExecutionMode) {
     // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
-    __SYNCTHREADS();
+    __kmpc_impl_syncthreads();
     if (GetThreadIdInBlock() == 0) {
       omptarget_nvptx_simpleMemoryManager.Release();
     }
     return;
   }
   __threadfence();
   ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   omptarget_nvptx_simpleMemoryManager.Release();
 }
 

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

 #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)
-#endif
-
 // arguments needed for L0 parallelism only.
 class omptarget_nvptx_SharedArgs {
 public:
   // All these methods must be called by the master thread only.
   INLINE void Init() {
     args  = buffer;
     nArgs = MAX_SHARED_ARGS;
   }

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

 //===--- omptarget-nvptx.cu - NVPTX OpenMP GPU initialization ---- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the initialization code for the GPU
 //
 //===----------------------------------------------------------------------===//
 
 #include "omptarget-nvptx.h"
+#include "target_impl.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // global data tables
 ////////////////////////////////////////////////////////////////////////////////
 
 extern __device__
     omptarget_nvptx_Queue<omptarget_nvptx_ThreadPrivateContext, OMP_STATE_COUNT>
         omptarget_nvptx_device_State[MAX_SM];
if (threadId == 0) {
     parallelLevel[0] =
         1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
   } else if (GetLaneId() == 0) {
     parallelLevel[GetWarpId()] =
         1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
   }
   if (!RequiresOMPRuntime) {
     // Runtime is not required - exit.
-    __SYNCTHREADS();
+    __kmpc_impl_syncthreads();
     return;
   }
 
   //
   // Team Context Initialization.
   //
   // In SPMD mode there is no master thread so use any cuda thread for team
   // context initialization.
   if (threadId == 0) {
     // Get a state object from the queue.
     omptarget_nvptx_threadPrivateContext =
         omptarget_nvptx_device_State[usedSlotIdx].Dequeue();
 
     omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
     omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
     // init team context
     currTeamDescr.InitTeamDescr();
   }
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
-  __SYNCTHREADS();
+  __kmpc_impl_syncthreads();
 
   omptarget_nvptx_TeamDescr &currTeamDescr = getMyTeamDescriptor();
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
 
   //
   // Initialize task descr for each thread.
   //
   omptarget_nvptx_TaskDescr *newTaskDescr =

 }
 
 EXTERN void __kmpc_spmd_kernel_deinit_v2(int16_t RequiresOMPRuntime) {
   // We're not going to pop the task descr stack of each thread since
   // there are no more parallel regions in SPMD mode.
   if (!RequiresOMPRuntime)
     return;
 
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
-  __SYNCTHREADS();
+  __kmpc_impl_syncthreads();
   int threadId = GetThreadIdInBlock();
   if (threadId == 0) {
     // Enqueue omp state object for use by another team.
     int slot = usedSlotIdx;
     omptarget_nvptx_device_State[slot].Enqueue(
         omptarget_nvptx_threadPrivateContext);
   }
 }
 
 // Return true if the current target region is executed in SPMD mode.
 EXTERN int8_t __kmpc_is_spmd_exec_mode() {
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_is_spmd_exec_mode\n");
   return isSPMDMode();
 }

openmp/libomptarget/deviceRTLs/nvptx/src/sync.cu

if (checkRuntimeUninitialized(loc_ref)) {
     PRINT0(LD_SYNC, "completed kmpc_barrier\n");
   }
 }
 
 // Emit a simple barrier call in SPMD mode.  Assumes the caller is in an L0
 // parallel region and that all worker threads participate.
 EXTERN void __kmpc_barrier_simple_spmd(kmp_Ident *loc_ref, int32_t tid) {
   PRINT0(LD_SYNC, "call kmpc_barrier_simple_spmd\n");
   // FIXME: use __syncthreads instead when the function copy is fixed in LLVM.
-  __SYNCTHREADS();
+  __kmpc_impl_syncthreads();
   PRINT0(LD_SYNC, "completed kmpc_barrier_simple_spmd\n");
 }
 
 // Emit a simple barrier call in Generic mode.  Assumes the caller is in an L0
 // parallel region and that all worker threads participate.
 EXTERN void __kmpc_barrier_simple_generic(kmp_Ident *loc_ref, int32_t tid) {
   int numberOfActiveOMPThreads = GetNumberOfThreadsInBlock() - WARPSIZE;
   // The #threads parameter must be rounded up to the WARPSIZE.

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); }
 
+#ifndef CUDA_VERSION
+#error CUDA_VERSION macro is undefined, something wrong with cuda.
+#endif
+
+INLINE void __kmpc_impl_syncthreads() {
+  // Use original __syncthreads if compiled by nvcc or clang >= 9.0.
+#if !defined(__clang__) || __clang_major__ >= 9
+  __syncthreads();
+#else
+  asm volatile("bar.sync %0;" : : "r"(0) : "memory");
+#endif
+}
+
 INLINE void __kmpc_impl_syncwarp(int32_t Mask) { __SYNCWARP(Mask); }
 
 #endif