[libomptarget][nfc] Move cuda threadfence functions behind kmpc_impl

[JonChesterfield]

Summary

[libomptarget][nfc] Move cuda threadfence functions behind kmpc_impl

Part of building code under common/ without requiring a cuda compiler

[ABataev]

LG

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

 
 INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   // we have protected the master warp from releasing from its barrier
   // due to a full workgroup barrier in the middle of a work function.
   // So it is ok to issue a full workgroup barrier here.
   __builtin_amdgcn_s_barrier();
 }
 
+EXTERN void __kmpc_impl_threadfence(void);
+EXTERN void __kmpc_impl_threadfence_block(void);
+EXTERN void __kmpc_impl_threadfence_system(void);
+
 // DEVICE versions of part of libc
 extern "C" {
 DEVICE __attribute__((noreturn)) void
 __assertfail(const char *, const char *, unsigned, const char *, size_t);
 INLINE static void __assert_fail(const char *__message, const char *__file,
                                  unsigned int __line, const char *__function) {
   __assertfail(__message, __file, __line, __function, sizeof(char));
 }
 DEVICE int printf(const char *, ...);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/common/src/loop.cu

if (schedule == kmp_sched_static_chunk) {
       if (chunk < 1)
         chunk = 1;
       omptarget_nvptx_threadPrivateContext->Chunk(tid) = chunk;
       omptarget_nvptx_threadPrivateContext->LoopUpperBound(tid) = ub;
       omptarget_nvptx_threadPrivateContext->NextLowerBound(tid) = lb;
       __kmpc_barrier(loc, threadId);
       if (tid == 0) {
         omptarget_nvptx_threadPrivateContext->Cnt() = 0;
-        __threadfence_block();
+        __kmpc_impl_threadfence_block();
       }
       __kmpc_barrier(loc, threadId);
       PRINT(LD_LOOP,
             "dispatch init (dyn) : num threads = %d, lb = %llu, ub = %" PRId64
             ", chunk %" PRIu64 "\n",
             (int)tnum,
             (unsigned long long)
                 omptarget_nvptx_threadPrivateContext->NextLowerBound(tid),

openmp/libomptarget/deviceRTLs/common/src/support.cu

 
 DEVICE void IncParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
   __kmpc_impl_syncwarp(Mask);
   __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt();
   unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] +=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
-    __threadfence();
+    __kmpc_impl_threadfence();
   }
   __kmpc_impl_syncwarp(Mask);
 }
 
 DEVICE void DecParallelLevel(bool ActiveParallel, __kmpc_impl_lanemask_t Mask) {
   __kmpc_impl_syncwarp(Mask);
   __kmpc_impl_lanemask_t LaneMaskLt = __kmpc_impl_lanemask_lt();
   unsigned Rank = __kmpc_impl_popc(Mask & LaneMaskLt);
   if (Rank == 0) {
     parallelLevel[GetWarpId()] -=
         (1 + (ActiveParallel ? OMP_ACTIVE_PARALLEL_LEVEL : 0));
-    __threadfence();
+    __kmpc_impl_threadfence();
   }
   __kmpc_impl_syncwarp(Mask);
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // get OpenMP number of procs
 
 // Get the number of processors in the device.

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

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Flush
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN void __kmpc_flush(kmp_Ident *loc) {
   PRINT0(LD_IO, "call kmpc_flush\n");
-  __threadfence();
+  __kmpc_impl_threadfence();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Vote
 ////////////////////////////////////////////////////////////////////////////////
 
 EXTERN __kmpc_impl_lanemask_t __kmpc_warp_active_thread_mask() {
   PRINT0(LD_IO, "call __kmpc_warp_active_thread_mask\n");

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

if (CurrentEndAddress < RequiredEndAddress) {
       }
 
       FrameP = StackP;
       StackP = (void *)RequiredEndAddress;
     }
   }
 
   // FIXME: Need to see the impact of doing it here.
-  __threadfence_block();
+  __kmpc_impl_threadfence_block();
 
   DSPRINT0(DSFLAG, "Exiting __kmpc_data_sharing_environment_begin\n");
 
   // All the threads in this warp get the frame they should work with.
   return FrameP;
 }
 
 EXTERN void __kmpc_data_sharing_environment_end(
if (!ActiveT) {
               (unsigned long long)StackP);
       DSPRINT(DSFLAG, "Restored frame ptr at: %016llx \n",
               (unsigned long long)FrameP);
       DSPRINT(DSFLAG, "Active threads: %08x \n", (unsigned)ActiveT);
     }
   }
 
   // FIXME: Need to see the impact of doing it here.
-  __threadfence_block();
+  __kmpc_impl_threadfence_block();
 
   DSPRINT0(DSFLAG, "Exiting __kmpc_data_sharing_environment_end\n");
   return;
 }
 
 EXTERN void *
 __kmpc_get_data_sharing_environment_frame(int32_t SourceThreadID,
                                           int16_t IsOMPRuntimeInitialized) {

 EXTERN void __kmpc_data_sharing_init_stack_spmd() {
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   // This function initializes the stack pointer with the pointer to the
   // statically allocated shared memory slots. The size of a shared memory
   // slot is pre-determined to be 256 bytes.
   if (GetThreadIdInBlock() == 0)
     data_sharing_init_stack_common();
 
-  __threadfence_block();
+  __kmpc_impl_threadfence_block();
 }
 
 INLINE static void* data_sharing_push_stack_common(size_t PushSize) {
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Expected initialized runtime.");
 
   // Only warp active master threads manage the stack.
   bool IsWarpMaster = (GetThreadIdInBlock() % WARPSIZE) == 0;
 

 // Pop the stack and free any memory which can be reclaimed.
 //
 // When the pop operation removes the last global memory slot,
 // reclaim all outstanding global memory slots since it is
 // likely we have reached the end of the kernel.
 EXTERN void __kmpc_data_sharing_pop_stack(void *FrameStart) {
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Expected initialized runtime.");
 
-  __threadfence_block();
+  __kmpc_impl_threadfence_block();
 
   if (GetThreadIdInBlock() % WARPSIZE == 0) {
     unsigned WID = GetWarpId();
 
     // Current slot
     __kmpc_data_sharing_slot *&SlotP = DataSharingState.SlotPtr[WID];
 
     // Pointer to next available stack.
if (GetThreadIdInBlock() == 0) {
       *frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
     }
     __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();
+  __kmpc_impl_threadfence();
 }
 
 EXTERN void __kmpc_restore_team_static_memory(int16_t isSPMDExecutionMode,
                                               int16_t is_shared) {
   if (is_shared)
     return;
   if (isSPMDExecutionMode) {
     __kmpc_impl_syncthreads();
     if (GetThreadIdInBlock() == 0) {
       omptarget_nvptx_simpleMemoryManager.Release();
     }
     return;
   }
-  __threadfence();
+  __kmpc_impl_threadfence();
   ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   omptarget_nvptx_simpleMemoryManager.Release();
 }
 

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

static int32_t nvptx_teams_reduce_nowait(int32_t global_tid, int32_t num_vars,
   SHARED volatile bool IsLastTeam;
 
   // Team masters of all teams write to the scratchpad.
   if (ThreadId == 0) {
     unsigned int *timestamp = GetTeamsReductionTimestamp();
     char *scratchpad = GetTeamsReductionScratchpad();
 
     scratchFct(reduce_data, scratchpad, TeamId, NumTeams);
-    __threadfence();
+    __kmpc_impl_threadfence();
 
     // atomicInc increments 'timestamp' and has a range [0, NumTeams-1].
     // It resets 'timestamp' back to 0 once the last team increments
     // this counter.
     unsigned val = atomicInc(timestamp, NumTeams - 1);
     IsLastTeam = val == NumTeams - 1;
   }
 
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_simple(kmp_Ident *loc,
   while (atomicCAS((uint32_t *)crit, 0, 1))
     ;
   return 1;
 }
 
 EXTERN void
 __kmpc_nvptx_teams_end_reduce_nowait_simple(kmp_Ident *loc, int32_t global_tid,
                                             kmp_CriticalName *crit) {
-  __threadfence_system();
+  __kmpc_impl_threadfence_system();
   (void)atomicExch((uint32_t *)crit, 0);
 }
 
 INLINE static bool isMaster(kmp_Ident *loc, uint32_t ThreadId) {
   return checkGenericMode(loc) || IsTeamMaster(ThreadId);
 }
 
 INLINE static uint32_t roundToWarpsize(uint32_t s) {
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   }
 
   if (IsMaster) {
     int ModBockId = TeamId % num_of_records;
     if (TeamId < num_of_records)
       lgcpyFct(global_buffer, ModBockId, reduce_data);
     else
       lgredFct(global_buffer, ModBockId, reduce_data);
-    __threadfence_system();
+    __kmpc_impl_threadfence_system();
 
     // Increment team counter.
     // This counter is incremented by all teams in the current
     // BUFFER_SIZE chunk.
     ChunkTeamCount = atomicInc((uint32_t *)&Cnt, num_of_records - 1);
   }
   // Synchronize
   if (checkSPMDMode(loc))

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

 
 INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   asm volatile("bar.sync %0, %1;"
                :
                : "r"(barrier), "r"(num_threads)
                : "memory");
 }
 
+INLINE void __kmpc_impl_threadfence(void) { __threadfence(); }
+INLINE void __kmpc_impl_threadfence_block(void) { __threadfence_block(); }
+INLINE void __kmpc_impl_threadfence_system(void) { __threadfence_system(); }
+
 #endif