[OpenMP] Expose libomptarget function to get HW thread id

[ggeorgakoudis]

Summary

The patch exposes the libomptarget runtime function that gets the hardware thread id through the kmpc API. This is to be used in SPMDization for checking the thread id to execute regions by a single thread in a block.

[jdoerfert]

Clang format, otherwise OK.

[ggeorgakoudis]

Fix formatting

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 
 
 
 __OMP_RTL(__kmpc_barrier, false, Void, IdentPtr, Int32)
 __OMP_RTL(__kmpc_cancel, false, Int32, IdentPtr, Int32, Int32)
 __OMP_RTL(__kmpc_cancel_barrier, false, Int32, IdentPtr, Int32)
 __OMP_RTL(__kmpc_flush, false, Void, IdentPtr)
 __OMP_RTL(__kmpc_global_thread_num, false, Int32, IdentPtr)
+__OMP_RTL(__kmpc_get_hardware_thread_id_in_block, false, Int32, )
 __OMP_RTL(__kmpc_fork_call, true, Void, IdentPtr, Int32, ParallelTaskPtr)
 __OMP_RTL(__kmpc_omp_taskwait, false, Int32, IdentPtr, Int32)
 __OMP_RTL(__kmpc_omp_taskyield, false, Int32, IdentPtr, Int32, /* Int */ Int32)
 __OMP_RTL(__kmpc_push_num_threads, false, Void, IdentPtr, Int32,
           /* Int */ Int32)
 __OMP_RTL(__kmpc_push_proc_bind, false, Void, IdentPtr, Int32, /* Int */ Int32)
 __OMP_RTL(__kmpc_omp_reg_task_with_affinity, false, Int32, IdentPtr, Int32,
           /* kmp_task_t */ VoidPtr, Int32,

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

 }
 
 EXTERN int GetNumberOfThreadsInBlock() {
   return get_workgroup_dim(__builtin_amdgcn_workgroup_id_x(),
                            __builtin_amdgcn_grid_size_x(),
                            __builtin_amdgcn_workgroup_size_x());
 }
 
-EXTERN unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+EXTERN unsigned GetWarpId() { return __kmpc_get_hardware_thread_id_in_block() / WARPSIZE; }
 EXTERN unsigned GetWarpSize() { return WARPSIZE; }
 EXTERN unsigned GetLaneId() {
   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
 }
 
 EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads() {
   return GetNumberOfThreadsInBlock();
 }
EXTERN void __kmpc_impl_threadfence_block() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
 }
 
 EXTERN void __kmpc_impl_threadfence_system() {
   __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
 }
 
 // Calls to the AMDGCN layer (assuming 1D layout)
-EXTERN int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
+EXTERN int __kmpc_get_hardware_thread_id_in_block() { return __builtin_amdgcn_workitem_id_x(); }
 EXTERN int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
 
 #pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/debug.h

 //   WARNING(flag, str, args)      : if warning flag is on, print the warning
 //   WON(flag)                     : return true if warning flag is on
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_NVPTX_DEBUG_H_
 #define _OMPTARGET_NVPTX_DEBUG_H_
 
 #include "common/device_environment.h"
 #include "target_interface.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // set desired level of debugging
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LD_SET_NONE 0ULL /* none */
 #define LD_SET_ALL -1ULL /* all */

 // implementation for debug
 ////////////////////////////////////////////////////////////////////////////////
 
 #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
 #include "common/support.h"
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
-  printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
-         (int)GetWarpId(), (int)GetLaneId(), parameters...);
+  printf(fmt, (int)GetBlockIdInKernel(),
+         (int)__kmpc_get_hardware_thread_id_in_block(), (int)GetWarpId(),
+         (int)GetLaneId(), parameters...);
 }
 
 #endif
 #if OMPTARGET_NVPTX_TEST
 
 template <typename... Arguments>
 NOINLINE static void check(bool cond, const char *fmt,
                            Arguments... parameters) {
   if (!cond) {
-    printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
-           (int)GetWarpId(), (int)GetLaneId(), parameters...);
+    printf(fmt, (int)GetBlockIdInKernel(),
+           (int)__kmpc_get_hardware_thread_id_in_block(), (int)GetWarpId(),
+           (int)GetLaneId(), parameters...);
     __builtin_trap();
   }
 }
 
 NOINLINE static void check(bool cond) {
   if (!cond)
     __builtin_trap();
 }

openmp/libomptarget/deviceRTLs/common/omptargeti.h

 // and functions.
 //
 //===----------------------------------------------------------------------===//
 
 ////////////////////////////////////////////////////////////////////////////////
 // Task Descriptor
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE omp_sched_t omptarget_nvptx_TaskDescr::GetRuntimeSched() const {
   // sched starts from 1..4; encode it as 0..3; so add 1 here
   uint8_t rc = (items.flags & TaskDescr_SchedMask) + 1;
   return (omp_sched_t)rc;
 }
 
 INLINE void omptarget_nvptx_TaskDescr::SetRuntimeSched(omp_sched_t sched) {
   // sched starts from 1..4; encode it as 0..3; so sub 1 here
   uint8_t val = ((uint8_t)sched) - 1;
INLINE void omptarget_nvptx_TaskDescr::InitLevelOneTaskDescr(
   // slow method
   // flag:
   //   default sched is static,
   //   dyn is off (unused now anyway, but may need to sample from host ?)
   //   in L1 parallel
 
   items.flags = TaskDescr_InPar | TaskDescr_IsParConstr; // set flag to parallel
   items.threadId =
-      GetThreadIdInBlock();   // get ids from cuda (only called for 1st level)
+      __kmpc_get_hardware_thread_id_in_block(); // get ids from cuda (only
+                                                // called for 1st level)
   items.runtimeChunkSize = 1; // preferred chunking statik with chunk 1
   prev = parentTaskDescr;
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyData(
     omptarget_nvptx_TaskDescr *sourceTaskDescr) {
   items = sourceTaskDescr->items;
 }

 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyFromWorkDescr(
     omptarget_nvptx_TaskDescr *workTaskDescr) {
   Copy(workTaskDescr);
   //
   // overwrite specific items;
   //
-  // The threadID should be GetThreadIdInBlock() % GetMasterThreadID().
-  // This is so that the serial master (first lane in the master warp)
-  // gets a threadId of 0.
-  // However, we know that this function is always called in a parallel
-  // region where only workers are active.  The serial master thread
-  // never enters this region.  When a parallel region is executed serially,
-  // the threadId is set to 0 elsewhere and the kmpc_serialized_* functions
-  // are called, which never activate this region.
+  // The threadID should be __kmpc_get_hardware_thread_id_in_block() %
+  // GetMasterThreadID(). This is so that the serial master (first lane in the
+  // master warp) gets a threadId of 0. However, we know that this function is
+  // always called in a parallel region where only workers are active.  The
+  // serial master thread never enters this region.  When a parallel region is
+  // executed serially, the threadId is set to 0 elsewhere and the
+  // kmpc_serialized_* functions are called, which never activate this region.
   items.threadId =
-      GetThreadIdInBlock(); // get ids from cuda (only called for 1st level)
+      __kmpc_get_hardware_thread_id_in_block(); // get ids from cuda (only
+                                                // called for 1st level)
 }
 
 INLINE void omptarget_nvptx_TaskDescr::CopyConvergentParent(
     omptarget_nvptx_TaskDescr *parentTaskDescr, uint16_t tid, uint16_t tnum) {
   CopyParent(parentTaskDescr);
   items.flags |= TaskDescr_InParL2P; // In L2+ parallelism
   items.threadId = tid;
 }

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

 #pragma omp allocate(WorkerSharedStack) allocator(omp_pteam_mem_alloc)
 
 template <typename AllocTy>
 static void *__kmpc_alloc_for_warp(AllocTy Alloc, unsigned Bytes,
                                    unsigned WarpBytes) {
   void *Ptr;
   __kmpc_impl_lanemask_t CurActive = __kmpc_impl_activemask();
   unsigned LeaderID = __kmpc_impl_ffs(CurActive) - 1;
-  bool IsWarpLeader = (GetThreadIdInBlock() % WARPSIZE) == LeaderID;
+  bool IsWarpLeader = (__kmpc_get_hardware_thread_id_in_block() % WARPSIZE) == LeaderID;
   if (IsWarpLeader)
     Ptr = Alloc();
   // Get address from the first active lane.
   int *FP = (int *)&Ptr;
   FP[0] = __kmpc_impl_shfl_sync(CurActive, FP[0], LeaderID);
   if (sizeof(Ptr) == 8)
     FP[1] = __kmpc_impl_shfl_sync(CurActive, FP[1], LeaderID);
   return (void *)&((char *)(Ptr))[(GetLaneId() - LeaderID) * Bytes];
 }
 
 EXTERN void *__kmpc_alloc_shared(size_t Bytes) {
   Bytes = Bytes + (Bytes % MinBytes);
-  int TID = GetThreadIdInBlock();
+  int TID = __kmpc_get_hardware_thread_id_in_block();
   if (__kmpc_is_generic_main_thread(TID)) {
     // Main thread alone, use shared memory if space is available.
     if (MainSharedStack.Usage[0] + Bytes <= MainSharedStack.MaxSize) {
       void *Ptr = &MainSharedStack.Data[MainSharedStack.Usage[0]];
       MainSharedStack.Usage[0] += Bytes;
       MainSharedStack.Sizes[MainSharedStack.SizeUsage[0]++] = Bytes;
       return Ptr;
     }
auto AllocGlobal = [&] {
     return SafeMalloc(WarpBytes, "AllocGlobalFallback");
   };
   return __kmpc_alloc_for_warp(AllocGlobal, Bytes, WarpBytes);
 }
 
 EXTERN void __kmpc_free_shared(void *Ptr) {
   __kmpc_impl_lanemask_t CurActive = __kmpc_impl_activemask();
   unsigned LeaderID = __kmpc_impl_ffs(CurActive) - 1;
-  bool IsWarpLeader = (GetThreadIdInBlock() % WARPSIZE) == LeaderID;
+  bool IsWarpLeader = (__kmpc_get_hardware_thread_id_in_block() % WARPSIZE) == LeaderID;
   __kmpc_syncwarp(CurActive);
   if (IsWarpLeader) {
     if (Ptr >= &MainSharedStack.Data[0] &&
         Ptr < &MainSharedStack.Data[MainSharedStack.MaxSize]) {
       unsigned Bytes = MainSharedStack.Sizes[--MainSharedStack.SizeUsage[0]];
       MainSharedStack.Usage[0] -= Bytes;
       return;
     }
EXTERN void __kmpc_get_team_static_memory(int16_t isSPMDExecutionMode,
                                           const void *buf, size_t size,
                                           int16_t is_shared,
                                           const void **frame) {
   if (is_shared) {
     *frame = buf;
     return;
   }
   if (isSPMDExecutionMode) {
-    if (GetThreadIdInBlock() == 0) {
+    if (__kmpc_get_hardware_thread_id_in_block() == 0) {
       *frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
     }
     __kmpc_impl_syncthreads();
     return;
   }
-  ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
+  ASSERT0(LT_FUSSY, __kmpc_get_hardware_thread_id_in_block() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   *frame = omptarget_nvptx_simpleMemoryManager.Acquire(buf, size);
   __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) {
+    if (__kmpc_get_hardware_thread_id_in_block() == 0) {
       omptarget_nvptx_simpleMemoryManager.Release();
     }
     return;
   }
   __kmpc_impl_threadfence();
-  ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
+  ASSERT0(LT_FUSSY, __kmpc_get_hardware_thread_id_in_block() == GetMasterThreadID(),
           "Must be called only in the target master thread.");
   omptarget_nvptx_simpleMemoryManager.Release();
 }
 
 #pragma omp end declare target

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

 EXTERN int omp_get_active_level(void) {
   int level = parallelLevel[GetWarpId()] > OMP_ACTIVE_PARALLEL_LEVEL ? 1 : 0;
   PRINT(LD_IO, "call omp_get_active_level() returns %d\n", level)
   return level;
 }
 
 EXTERN int omp_get_ancestor_thread_num(int level) {
   if (__kmpc_is_spmd_exec_mode())
-    return level == 1 ? GetThreadIdInBlock() : 0;
+    return level == 1 ? __kmpc_get_hardware_thread_id_in_block() : 0;
   int rc = -1;
   // If level is 0 or all parallel regions are not active - return 0.
   unsigned parLevel = parallelLevel[GetWarpId()];
   if (level == 1 && parLevel > OMP_ACTIVE_PARALLEL_LEVEL) {
     int totLevel = omp_get_level();
     if (level <= totLevel) {
       omptarget_nvptx_TaskDescr *currTaskDescr =
           getMyTopTaskDescriptor(/*isSPMDExecutionMode=*/false);

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

 
 static void __kmpc_generic_kernel_init() {
   PRINT(LD_IO, "call to __kmpc_kernel_init with version %f\n",
         OMPTARGET_NVPTX_VERSION);
 
   if (GetLaneId() == 0)
     parallelLevel[GetWarpId()] = 0;
 
-  int threadIdInBlock = GetThreadIdInBlock();
+  int threadIdInBlock = __kmpc_get_hardware_thread_id_in_block();
   if (threadIdInBlock != GetMasterThreadID())
     return;
 
   setExecutionParameters(Generic, RuntimeInitialized);
   ASSERT0(LT_FUSSY, threadIdInBlock == GetMasterThreadID(),
           "__kmpc_kernel_init() must be called by team master warp only!");
   PRINT0(LD_IO, "call to __kmpc_kernel_init for master\n");
 
static void __kmpc_generic_kernel_deinit() {
   omptarget_nvptx_workFn = 0;
 }
 
 static void __kmpc_spmd_kernel_init(bool RequiresFullRuntime) {
   PRINT0(LD_IO, "call to __kmpc_spmd_kernel_init\n");
 
   setExecutionParameters(Spmd, RequiresFullRuntime ? RuntimeInitialized
                          : RuntimeUninitialized);
-  int threadId = GetThreadIdInBlock();
+  int threadId = __kmpc_get_hardware_thread_id_in_block();
   if (threadId == 0) {
     usedSlotIdx = __kmpc_impl_smid() % MAX_SM;
   }
 
   if (GetLaneId() == 0) {
     parallelLevel[GetWarpId()] =
         1 + (GetNumberOfThreadsInBlock() > 1 ? OMP_ACTIVE_PARALLEL_LEVEL : 0);
   }

 
 static void __kmpc_spmd_kernel_deinit(bool RequiresFullRuntime) {
   // We're not going to pop the task descr stack of each thread since
   // there are no more parallel regions in SPMD mode.
   if (!RequiresFullRuntime)
     return;
 
   __kmpc_impl_syncthreads();
-  int threadId = GetThreadIdInBlock();
+  int threadId = __kmpc_get_hardware_thread_id_in_block();
   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() {
   return (execution_param & ModeMask) == Spmd;
 }
 
 EXTERN int8_t __kmpc_is_generic_main_thread(kmp_int32 Tid) {
   return !__kmpc_is_spmd_exec_mode() && GetMasterThreadID() == Tid;
 }
 
 EXTERN bool __kmpc_kernel_parallel(void**WorkFn);
 
 static void __kmpc_target_region_state_machine(ident_t *Ident) {
 
-  int TId = GetThreadIdInBlock();
+  int TId = __kmpc_get_hardware_thread_id_in_block();
   do {
     void* WorkFn = 0;
 
     // Wait for the signal that we have a new work function.
     __kmpc_barrier_simple_spmd(Ident, TId);
 
 
     // Retrieve the work function from the runtime.
static void __kmpc_target_region_state_machine(ident_t *Ident) {
 
   } while (true);
 }
 
 EXTERN
 int32_t __kmpc_target_init(ident_t *Ident, bool IsSPMD,
                            bool UseGenericStateMachine,
                            bool RequiresFullRuntime) {
-  int TId = GetThreadIdInBlock();
+  int TId = __kmpc_get_hardware_thread_id_in_block();
   if (IsSPMD)
     __kmpc_spmd_kernel_init(RequiresFullRuntime);
   else
     __kmpc_generic_kernel_init();
 
    if (IsSPMD) {
     __kmpc_barrier_simple_spmd(Ident, TId);
      return -1;

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

EXTERN void __kmpc_kernel_prepare_parallel(void *WorkFn) {
 
   if (NumThreadsClause != 0) {
     // Reset request to avoid propagating to successive #parallel
     NumThreadsClause = 0;
   }
 
   ASSERT(LT_FUSSY, NumThreads > 0, "bad thread request of %d threads",
          (int)NumThreads);
-  ASSERT0(LT_FUSSY, GetThreadIdInBlock() == GetMasterThreadID(),
+  ASSERT0(LT_FUSSY, __kmpc_get_hardware_thread_id_in_block() == GetMasterThreadID(),
           "only team master can create parallel");
 
   // Set number of threads on work descriptor.
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
   workDescr.WorkTaskDescr()->CopyToWorkDescr(currTaskDescr);
   threadsInTeam = NumThreads;
 }
 
EXTERN bool __kmpc_kernel_parallel(void **WorkFn) {
   // If this is the termination signal from the master, quit early.
   if (!*WorkFn) {
     PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_parallel finished\n");
     return false;
   }
 
   // Only the worker threads call this routine and the master warp
   // never arrives here.  Therefore, use the nvptx thread id.
-  int threadId = GetThreadIdInBlock();
+  int threadId = __kmpc_get_hardware_thread_id_in_block();
   omptarget_nvptx_WorkDescr &workDescr = getMyWorkDescriptor();
   // Set to true for workers participating in the parallel region.
   bool isActive = false;
   // Initialize state for active threads.
   if (threadId < threadsInTeam) {
     // init work descriptor from workdesccr
     omptarget_nvptx_TaskDescr *newTaskDescr =
         omptarget_nvptx_threadPrivateContext->Level1TaskDescr(threadId);

 
 EXTERN void __kmpc_kernel_end_parallel() {
   // pop stack
   PRINT0(LD_IO | LD_PAR, "call to __kmpc_kernel_end_parallel\n");
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Expected initialized runtime.");
 
   // Only the worker threads call this routine and the master warp
   // never arrives here.  Therefore, use the nvptx thread id.
-  int threadId = GetThreadIdInBlock();
+  int threadId = __kmpc_get_hardware_thread_id_in_block();
   omptarget_nvptx_TaskDescr *currTaskDescr = getMyTopTaskDescriptor(threadId);
   omptarget_nvptx_threadPrivateContext->SetTopLevelTaskDescr(
       threadId, currTaskDescr->GetPrevTaskDescr());
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // support for parallel that goes sequential
 ////////////////////////////////////////////////////////////////////////////////

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

while (mask > 0) {
     mask = curr_size / 2;
   }
 }
 
 #if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 700
 INLINE static uint32_t
 gpu_irregular_simd_reduce(void *reduce_data, kmp_ShuffleReductFctPtr shflFct) {
   uint32_t size, remote_id, physical_lane_id;
-  physical_lane_id = GetThreadIdInBlock() % WARPSIZE;
+  physical_lane_id = __kmpc_get_hardware_thread_id_in_block() % WARPSIZE;
   __kmpc_impl_lanemask_t lanemask_lt = __kmpc_impl_lanemask_lt();
   __kmpc_impl_lanemask_t Liveness = __kmpc_impl_activemask();
   uint32_t logical_lane_id = __kmpc_impl_popc(Liveness & lanemask_lt) * 2;
   __kmpc_impl_lanemask_t lanemask_gt = __kmpc_impl_lanemask_gt();
   do {
     Liveness = __kmpc_impl_activemask();
     remote_id = __kmpc_impl_ffs(Liveness & lanemask_gt);
     size = __kmpc_impl_popc(Liveness);
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
 
   // Volta execution model:
   // For the Generic execution mode a parallel region either has 1 thread and
   // beyond that, always a multiple of 32. For the SPMD execution mode we may
   // have any number of threads.
   if ((NumThreads % WARPSIZE == 0) || (WarpId < WarpsNeeded - 1))
     gpu_regular_warp_reduce(reduce_data, shflFct);
   else if (NumThreads > 1) // Only SPMD execution mode comes thru this case.
     gpu_irregular_warp_reduce(reduce_data, shflFct,
                               /*LaneCount=*/NumThreads % WARPSIZE,
-                              /*LaneId=*/GetThreadIdInBlock() % WARPSIZE);
+                              /*LaneId=*/__kmpc_get_hardware_thread_id_in_block() % WARPSIZE);
 
   // When we have more than [warpsize] number of threads
   // a block reduction is performed here.
   //
   // Only L1 parallel region can enter this if condition.
   if (NumThreads > WARPSIZE) {
     // Gather all the reduced values from each warp
     // to the first warp.
     cpyFct(reduce_data, WarpsNeeded);
 
     if (WarpId == 0)
       gpu_irregular_warp_reduce(reduce_data, shflFct, WarpsNeeded,
                                 BlockThreadId);
   }
   return BlockThreadId == 0;
 #else
   __kmpc_impl_lanemask_t Liveness = __kmpc_impl_activemask();
   if (Liveness == __kmpc_impl_all_lanes) // Full warp
     gpu_regular_warp_reduce(reduce_data, shflFct);
   else if (!(Liveness & (Liveness + 1))) // Partial warp but contiguous lanes
     gpu_irregular_warp_reduce(reduce_data, shflFct,
                               /*LaneCount=*/__kmpc_impl_popc(Liveness),
-                              /*LaneId=*/GetThreadIdInBlock() % WARPSIZE);
+                              /*LaneId=*/__kmpc_get_hardware_thread_id_in_block() % WARPSIZE);
   else if (!isRuntimeUninitialized) // Dispersed lanes. Only threads in L2
                                     // parallel region may enter here; return
                                     // early.
     return gpu_irregular_simd_reduce(reduce_data, shflFct);
 
   // When we have more than [warpsize] number of threads
   // a block reduction is performed here.
   //
EXTERN int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
     kmp_Ident *loc, int32_t global_tid, void *global_buffer,
     int32_t num_of_records, void *reduce_data, kmp_ShuffleReductFctPtr shflFct,
     kmp_InterWarpCopyFctPtr cpyFct, kmp_ListGlobalFctPtr lgcpyFct,
     kmp_ListGlobalFctPtr lgredFct, kmp_ListGlobalFctPtr glcpyFct,
     kmp_ListGlobalFctPtr glredFct) {
 
   // Terminate all threads in non-SPMD mode except for the master thread.
   if (!__kmpc_is_spmd_exec_mode() &&
-      !__kmpc_is_generic_main_thread(GetThreadIdInBlock()))
+      !__kmpc_is_generic_main_thread(__kmpc_get_hardware_thread_id_in_block()))
     return 0;
 
   uint32_t ThreadId = GetLogicalThreadIdInBlock();
 
   // In non-generic mode all workers participate in the teams reduction.
   // In generic mode only the team master participates in the teams
   // reduction because the workers are waiting for parallel work.
   uint32_t NumThreads =

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

 
 // This function may be called in a parallel region by the workers
 // or a serial region by the master.  If the master (whose CUDA thread
 // id is GetMasterThreadID()) calls this routine, we return 0 because
 // it is a shadow for the first worker.
 int GetLogicalThreadIdInBlock() {
   // Implemented using control flow (predication) instead of with a modulo
   // operation.
-  int tid = GetThreadIdInBlock();
+  int tid = __kmpc_get_hardware_thread_id_in_block();
   if (__kmpc_is_generic_main_thread(tid))
     return 0;
   else
     return tid;
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // OpenMP Thread Support Layer
 //
 ////////////////////////////////////////////////////////////////////////////////
 
 int GetOmpThreadId() {
-  int tid = GetThreadIdInBlock();
+  int tid = __kmpc_get_hardware_thread_id_in_block();
   if (__kmpc_is_generic_main_thread(tid))
     return 0;
   // omp_thread_num
   int rc;
   if (__kmpc_parallel_level() > 1) {
     rc = 0;
   } else if (__kmpc_is_spmd_exec_mode()) {
     rc = tid;

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

 }
 
 EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   unsigned int Mask;
   asm volatile("activemask.b32 %0;" : "=r"(Mask));
   return Mask;
 }
 
-EXTERN void __kmpc_impl_syncthreads() { 
+EXTERN void __kmpc_impl_syncthreads() {
   int barrier = 2;
   asm volatile("barrier.sync %0;"
                :
                : "r"(barrier)
                : "memory");
 }
 
 EXTERN void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
asm volatile("barrier.sync %0, %1;"
                : "memory");
 }
 
 EXTERN void __kmpc_impl_threadfence() { __nvvm_membar_gl(); }
 EXTERN void __kmpc_impl_threadfence_block() { __nvvm_membar_cta(); }
 EXTERN void __kmpc_impl_threadfence_system() { __nvvm_membar_sys(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
-EXTERN int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
+EXTERN int __kmpc_get_hardware_thread_id_in_block() { return __nvvm_read_ptx_sreg_tid_x(); }
 EXTERN int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
 EXTERN int GetNumberOfBlocksInKernel() {
   return __nvvm_read_ptx_sreg_nctaid_x();
 }
 EXTERN int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
-EXTERN unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+EXTERN unsigned GetWarpId() { return __kmpc_get_hardware_thread_id_in_block() / WARPSIZE; }
 EXTERN unsigned GetWarpSize() { return WARPSIZE; }
-EXTERN unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
+EXTERN unsigned GetLaneId() { return __kmpc_get_hardware_thread_id_in_block() & (WARPSIZE - 1); }
 
 // Atomics
 uint32_t __kmpc_atomic_add(uint32_t *Address, uint32_t Val) {
   return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 uint32_t __kmpc_atomic_inc(uint32_t *Address, uint32_t Val) {
   return __nvvm_atom_inc_gen_ui(Address, Val);
 }

openmp/libomptarget/deviceRTLs/target_interface.h

 //===------------- target_interface.h - Target interfaces --------- C++ -*-===//
 //
 // 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 interfaces that must be implemented by each target.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_TARGET_INTERFACE_H_
 #define _OMPTARGET_TARGET_INTERFACE_H_
 
 #include "target_impl.h"
 
 // Calls to the NVPTX layer (assuming 1D layout)
-EXTERN int GetThreadIdInBlock();
+EXTERN int __kmpc_get_hardware_thread_id_in_block();
 EXTERN int GetBlockIdInKernel();
 EXTERN int GetNumberOfBlocksInKernel();
 EXTERN int GetNumberOfThreadsInBlock();
 EXTERN unsigned GetWarpId();
 EXTERN unsigned GetWarpSize();
 EXTERN unsigned GetLaneId();
 
 // Atomics