[libomptarget][nfc] Use cuda variable wrappers from support.h

[JonChesterfield]

Summary

[libomptarget][nfc] Use cuda variable wrappers from support.h
Reimplementation of D69693, after the revert of D69885

Use the wrappers in support.h for cuda builtin variables at all call sites.
Localises use of cuda and removes WARPSIZE==32 assumption in debug.h.

[jdoerfert]

LGTM. (I think pure NFC cleanups like this you can commit w/o review.)

[ABataev]

In D70186#1744100, @jdoerfert wrote:

LGTM. (I think pure NFC cleanups like this you can commit w/o review.)

Need to check that it does not break the build, give me some time.

[JonChesterfield]

Thanks. In general I agree re: NFC. At present I'm being cautious after the non-reproducible build errors of last week.

The missing symbols were root caused - __forceinline__ is not equal to always_inline - but duplicate symbols have me at a loss, especially when the supporti.h => support.cu patch didn't show the issue. Hopefully Alexey will run into them again on one of the subdivided patches.

[ABataev]

In D70186#1744191, @JonChesterfield wrote:

Thanks. In general I agree re: NFC. At present I'm being cautious after the non-reproducible build errors of last week.

The missing symbols were root caused - __forceinline__ is not equal to always_inline - but duplicate symbols have me at a loss, especially when the supporti.h => support.cu patch didn't show the issue. Hopefully Alexey will run into them again on one of the subdivided patches.

Tested, works fine.

[jdoerfert]

In D70186#1744200, @ABataev wrote:

In D70186#1744191, @JonChesterfield wrote:

Thanks. In general I agree re: NFC. At present I'm being cautious after the non-reproducible build errors of last week.

The missing symbols were root caused - __forceinline__ is not equal to always_inline - but duplicate symbols have me at a loss, especially when the supporti.h => support.cu patch didn't show the issue. Hopefully Alexey will run into them again on one of the subdivided patches.

Tested, works fine.

Thx!

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.
 INLINE static bool IsWarpMasterActiveThread() {
   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;
 }

 // the slot.
 EXTERN void
 __kmpc_initialize_data_sharing_environment(__kmpc_data_sharing_slot *rootS,
                                            size_t InitialDataSize) {
   ASSERT0(LT_FUSSY, isRuntimeInitialized(), "Runtime must be initialized.");
   DSPRINT0(DSFLAG_INIT,
            "Entering __kmpc_initialize_data_sharing_environment\n");
 
-  unsigned WID = getWarpId();
+  unsigned WID = GetWarpId();
   DSPRINT(DSFLAG_INIT, "Warp ID: %u\n", WID);
 
   omptarget_nvptx_TeamDescr *teamDescr =
       &omptarget_nvptx_threadPrivateContext->TeamContext();
   __kmpc_data_sharing_slot *RootS =
       teamDescr->RootS(WID, IsMasterThread(isSPMDMode()));
 
   DataSharingState.SlotPtr[WID] = RootS;
EXTERN void *__kmpc_data_sharing_environment_begin(
   // data sharing.
   if (!IsOMPRuntimeInitialized)
     return (void *)&DataSharingState;
 
   DSPRINT(DSFLAG, "Data Size %016llx\n", (unsigned long long)SharingDataSize);
   DSPRINT(DSFLAG, "Default Data Size %016llx\n",
           (unsigned long long)SharingDefaultDataSize);
 
-  unsigned WID = getWarpId();
+  unsigned WID = GetWarpId();
   __kmpc_impl_lanemask_t CurActiveThreads = __kmpc_impl_activemask();
 
   __kmpc_data_sharing_slot *&SlotP = DataSharingState.SlotPtr[WID];
   void *&StackP = DataSharingState.StackPtr[WID];
   void * volatile &FrameP = DataSharingState.FramePtr[WID];
   __kmpc_impl_lanemask_t &ActiveT = DataSharingState.ActiveThreads[WID];
 
   DSPRINT0(DSFLAG, "Save current slot/stack values.\n");

 
 EXTERN void __kmpc_data_sharing_environment_end(
     __kmpc_data_sharing_slot **SavedSharedSlot, void **SavedSharedStack,
     void **SavedSharedFrame, __kmpc_impl_lanemask_t *SavedActiveThreads,
     int32_t IsEntryPoint) {
 
   DSPRINT0(DSFLAG, "Entering __kmpc_data_sharing_environment_end\n");
 
-  unsigned WID = getWarpId();
+  unsigned WID = GetWarpId();
 
   if (IsEntryPoint) {
     if (IsWarpMasterActiveThread()) {
       DSPRINT0(DSFLAG, "Doing clean up\n");
 
       // The master thread cleans the saved slot, because this is an environment
       // only for the master.
       __kmpc_data_sharing_slot *S = IsMasterThread(isSPMDMode())

 // Initialize data sharing data structure. This function needs to be called
 // once at the beginning of a data sharing context (coincides with the kernel
 // initialization). This function is called in SPMD mode only.
 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 (threadIdx.x == 0)
+  if (GetThreadIdInBlock() == 0)
     data_sharing_init_stack_common();
 
   __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;
 
   // Add worst-case padding to DataSize so that future stack allocations are
   // correctly aligned.
   const size_t Alignment = 8;
   PushSize = (PushSize + (Alignment - 1)) / Alignment * Alignment;
 
   // Frame pointer must be visible to all workers in the same warp.
-  const unsigned WID = getWarpId();
+  const unsigned WID = GetWarpId();
   void *FrameP = 0;
   __kmpc_impl_lanemask_t CurActive = __kmpc_impl_activemask();
 
   if (IsWarpMaster) {
     // SlotP will point to either the shared memory slot or an existing
     // global memory slot.
     __kmpc_data_sharing_slot *&SlotP = DataSharingState.SlotPtr[WID];
     void *&StackP = DataSharingState.StackPtr[WID];
EXTERN void *__kmpc_data_sharing_push_stack(size_t DataSize,
   // TODO: change WARPSIZE to the number of active threads in the warp.
   size_t PushSize = (isRuntimeUninitialized() || IsMasterThread(isSPMDMode()))
                         ? DataSize
                         : WARPSIZE * DataSize;
 
   // Compute the start address of the frame of each thread in the warp.
   uintptr_t FrameStartAddress =
       (uintptr_t) data_sharing_push_stack_common(PushSize);
-  FrameStartAddress += (uintptr_t) (getLaneId() * DataSize);
+  FrameStartAddress += (uintptr_t) (GetLaneId() * DataSize);
   return (void *)FrameStartAddress;
 }
 
 // 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();
 
   if (GetThreadIdInBlock() % WARPSIZE == 0) {
-    unsigned WID = getWarpId();
+    unsigned WID = GetWarpId();
 
     // Current slot
     __kmpc_data_sharing_slot *&SlotP = DataSharingState.SlotPtr[WID];
 
     // Pointer to next available stack.
     void *&StackP = DataSharingState.StackPtr[WID];
 
     // Pop the frame.

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

 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // implemtation for debug
 ////////////////////////////////////////////////////////////////////////////////
 
 #if OMPTARGET_NVPTX_DEBUG || OMPTARGET_NVPTX_TEST || OMPTARGET_NVPTX_WARNING
 #include <stdio.h>
-#include "target_impl.h"
+#include "support.h"
 
 template <typename... Arguments>
 NOINLINE static void log(const char *fmt, Arguments... parameters) {
-  printf(fmt, (int)blockIdx.x, (int)threadIdx.x, (int)(threadIdx.x / WARPSIZE),
-         (int)(threadIdx.x & 0x1F), parameters...);
+  printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
+         (int)GetWarpId(), (int)GetLaneId(), parameters...);
 }
 
 #endif
 #if OMPTARGET_NVPTX_TEST
 #include <assert.h>
 
 template <typename... Arguments>
 NOINLINE static void check(bool cond, const char *fmt,
                            Arguments... parameters) {
   if (!cond)
-    printf(fmt, (int)blockIdx.x, (int)threadIdx.x,
-           (int)(threadIdx.x / WARPSIZE), (int)(threadIdx.x & 0x1F),
-           parameters...);
+    printf(fmt, (int)GetBlockIdInKernel(), (int)GetThreadIdInBlock(),
+           (int)GetWarpId(), (int)GetLaneId(), parameters...);
   assert(cond);
 }
 
 NOINLINE static void check(bool cond) { assert(cond); }
 #endif
 
 // set flags that are tested (inclusion properties)
 

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

EXTERN void omp_set_lock(omp_lock_t *lock) {
 
   // TODO: not sure spinning is a good idea here..
   while (atomicCAS(lock, UNSET, SET) != UNSET) {
     clock_t start = clock();
     clock_t now;
     for (;;) {
       now = clock();
       clock_t cycles = now > start ? now - start : now + (0xffffffff - start);
-      if (cycles >= __OMP_SPIN * blockIdx.x) {
+      if (cycles >= __OMP_SPIN * GetBlockIdInKernel()) {
         break;
       }
     }
   } // wait for 0 to be the read value
 
   PRINT0(LD_IO, "call omp_set_lock()\n");
 }
 

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

 DEVICE int GetThreadIdInBlock() { return threadIdx.x; }
 
 DEVICE int GetBlockIdInKernel() { return blockIdx.x; }
 
 DEVICE int GetNumberOfBlocksInKernel() { return gridDim.x; }
 
 DEVICE int GetNumberOfThreadsInBlock() { return blockDim.x; }
 
-DEVICE unsigned GetWarpId() { return threadIdx.x / WARPSIZE; }
+DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 
-DEVICE unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); }
+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.
 //      If NumThreads is 97, master id is 96.
 //      If NumThreads is 1024, master id is 992.
 //
 // Called in Generic Execution Mode only.
-DEVICE int GetMasterThreadID() { return (blockDim.x - 1) & ~(WARPSIZE - 1); }
+DEVICE int GetMasterThreadID() { return (GetNumberOfThreadsInBlock() - 1) & ~(WARPSIZE - 1); }
 
 // The last warp is reserved for the master; other warps are workers.
 // Called in Generic Execution Mode only.
 DEVICE int GetNumberOfWorkersInTeam() { return GetMasterThreadID(); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // get thread id in team