[libomptarget][nfc] Move GetWarp/LaneId functions into per arch code

[JonChesterfield]

Summary

[libomptarget][nfc] Move GetWarp/LaneId functions into per arch code

No code change for nvptx. Amdgcn currently has two implementations of GetLaneId,
this patch keeps the one a colleague considered to be superior for our ISA.

GetWarpId is currently the same function for amdgcn and nvptx, but I think it's
cleaner to keep it grouped with all the others than to keep it in support.cu.

[JonChesterfield]

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

17

^ this is the preferred implementation, which is unfortunately different to the generic one

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

194

^nvptx uses exactly the same functions as before

[jdoerfert]

LGTM

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

 DEVICE void __kmpc_impl_threadfence_block(void);
 DEVICE void __kmpc_impl_threadfence_system(void);
 
 // Calls to the AMDGCN layer (assuming 1D layout)
 INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
 INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
 DEVICE int GetNumberOfBlocksInKernel();
 DEVICE int GetNumberOfThreadsInBlock();
+DEVICE unsigned GetWarpId();
+DEVICE unsigned GetLaneId();
 
 DEVICE bool __kmpc_impl_is_first_active_thread();
 
 // Locks
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock);
 DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *lock);
 DEVICE void __kmpc_impl_set_lock(omp_lock_t *lock);
 DEVICE void __kmpc_impl_unset_lock(omp_lock_t *lock);

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

 //===------- target_impl.hip - AMDGCN OpenMP GPU implementation --- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 
 #include "target_impl.h"
 
 // Implementations initially derived from hcc
 
-static DEVICE uint32_t getLaneId(void) {
-  return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
-}
-
 // Initialized with a 64-bit mask with bits set in positions less than the
 // thread's lane number in the warp
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
-  uint32_t lane = getLaneId();
+  uint32_t lane = GetLaneId();
   int64_t ballot = __kmpc_impl_activemask();
   uint64_t mask = ((uint64_t)1 << lane) - (uint64_t)1;
   return mask & ballot;
 }
 
 // Initialized with a 64-bit mask with bits set in positions greater than the
 // thread's lane number in the warp
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
-  uint32_t lane = getLaneId();
+  uint32_t lane = GetLaneId();
   if (lane == (WARPSIZE - 1))
     return 0;
   uint64_t ballot = __kmpc_impl_activemask();
   uint64_t mask = (~((uint64_t)0)) << (lane + 1);
   return mask & ballot;
 }
 
 DEVICE double __kmpc_impl_get_wtick() { return ((double)1E-9); }
 
 EXTERN uint64_t __clock64();
 DEVICE double __kmpc_impl_get_wtime() {
   return ((double)1.0 / 745000000.0) * __clock64();
 }
 
 // Warp vote function
 DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   return __builtin_amdgcn_read_exec();
 }
 
 DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t var,
                                      int32_t srcLane) {
   int width = WARPSIZE;
-  int self = getLaneId();
+  int self = GetLaneId();
   int index = srcLane + (self & ~(width - 1));
   return __builtin_amdgcn_ds_bpermute(index << 2, var);
 }
 
 DEVICE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t var,
                                           uint32_t laneDelta, int32_t width) {
-  int self = getLaneId();
+  int self = GetLaneId();
   int index = self + laneDelta;
   index = (int)(laneDelta + (self & (width - 1))) >= width ? self : index;
   return __builtin_amdgcn_ds_bpermute(index << 2, var);
 }
 
 EXTERN uint64_t __ockl_get_local_size(uint32_t);
 EXTERN uint64_t __ockl_get_num_groups(uint32_t);
 DEVICE int GetNumberOfBlocksInKernel() { return __ockl_get_num_groups(0); }
 DEVICE int GetNumberOfThreadsInBlock() { return __ockl_get_local_size(0); }
+DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+DEVICE unsigned GetLaneId() {
+  return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
+}

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

 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // support: get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 //
-// Calls to the NVPTX layer  (assuming 1D layout)
-//
-////////////////////////////////////////////////////////////////////////////////
-
-DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
-
-DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
-
-////////////////////////////////////////////////////////////////////////////////
-//
 // Calls to the Generic Scheme Implementation Layer (assuming 1D layout)
 //
 ////////////////////////////////////////////////////////////////////////////////
 
 // The master thread id is the first thread (lane) of the last warp.
 // Thread id is 0 indexed.
 // E.g: If NumThreads is 33, master id is 32.
 //      If NumThreads is 64, master id is 32.

openmp/libomptarget/deviceRTLs/common/support.h

 DEVICE bool checkGenericMode(kmp_Ident *loc);
 DEVICE bool checkRuntimeUninitialized(kmp_Ident *loc);
 DEVICE bool checkRuntimeInitialized(kmp_Ident *loc);
 
 ////////////////////////////////////////////////////////////////////////////////
 // get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
-// get low level ids of resources
-DEVICE unsigned GetWarpId();
-DEVICE unsigned GetLaneId();
-
 // get global ids to locate tread/team info (constant regardless of OMP)
 DEVICE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
 DEVICE int GetMasterThreadID();
 DEVICE int GetNumberOfWorkersInTeam();
 
 // get OpenMP thread and team ids
 DEVICE int GetOmpThreadId(int threadId,
                           bool isSPMDExecutionMode);    // omp_thread_num

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

 INLINE void __kmpc_impl_threadfence_block(void) { __threadfence_block(); }
 INLINE void __kmpc_impl_threadfence_system(void) { __threadfence_system(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
 INLINE int GetThreadIdInBlock() { return threadIdx.x; }
 INLINE int GetBlockIdInKernel() { return blockIdx.x; }
 INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; }
 INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+INLINE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+INLINE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 
 // Return true if this is the first active thread in the warp.
 INLINE bool __kmpc_impl_is_first_active_thread() {
   unsigned long long Mask = __kmpc_impl_activemask();
   unsigned long long ShNum = WARPSIZE - (GetThreadIdInBlock() % WARPSIZE);
   unsigned long long Sh = Mask << ShNum;
   // Truncate Sh to the 32 lower bits
   return (unsigned)Sh == 0;
 }