[nfc][libomptarget] Warp size aware logging by using builtinvar abstraction consistently. Last part of D69423

[JonChesterfield]

Summary

[nfc][libomptarget] Warp size aware logging by using builtinvar abstraction consistently. Last part of D69423

Uses functions from support.h to derive values used in logging from WARPSIZE.

Also updates the rest of the deviceRTL to use this existing abstraction everywhere.

[JonChesterfield]

• Use functions from support.h instead of adding to target_impl

[JonChesterfield]

• Use cuda builtin_var abstraction everywhere

[JonChesterfield]

A subsequent patch will be proposed to move the six functions labelled 'Calls to the NVPTX layer' in support.h into target_impl.

That will allow me to delete the ugliest third of the api shim used by amdgcn to work around not having a cuda compiler. Code like:

#define __CUDA_DEVICE_BUILTIN(FIELD, INTRINSIC)                                \
  __declspec(property(get = __fetch_builtin_##FIELD)) unsigned int FIELD;      \
  static inline __attribute__((always_inline))                                 \
      __attribute__((device)) unsigned int __fetch_builtin_##FIELD(void) {     \
    return INTRINSIC;                                                          \
  }
//...
__CUDA_DEVICE_BUILTIN(x, __nvvm_read_ptx_sreg_tid_x());
//...
define i32 @__nvvm_read_ptx_sreg_tid_x() #2 {
  %id.i = tail call i32 @llvm.amdgcn.workitem.id.x() #3
  ret i32 %id.i
}

[JonChesterfield]

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

109

LLVM transforms the mod with constant to mask with constant - 1 at O1.

[ABataev]

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

126

Hmm, why here it is not ttansformed into -1?

[JonChesterfield]

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

126

The ~. Change only converts the variable access into a function call.

Codegens as (blockDim.x - 1) & -32 which seems reasonable.

[JonChesterfield]

There's some collateral from this. Support.h is missing include guards and isn't safe to include multiple times. If guards are introduced, the inline functions that are defined in supporti.h generate warnings from Wundefined-inline. If supporti.h is appended to support, it turns out to be missing a number of headers.

I'm going detour to fix up some header dependencies then recreate an equivalent diff to this one.

[JonChesterfield]

Fixing the headers is disproportionately expensive and will make a mess of the source dir. Lots of little headers.

Posted an alternate to openmp-dev suggesting we implement functions in source files instead, working around nvcc lacking lto with a unity build.

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 "support.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;
 }

 // 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.");
 

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/supporti.h

 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 threadIdx.x / WARPSIZE; }
+INLINE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 
-INLINE unsigned GetLaneId() { return threadIdx.x & (WARPSIZE - 1); }
+INLINE unsigned GetLaneId() { return GetThreadIdInBlock() % WARPSIZE; }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // 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.
-INLINE int GetMasterThreadID() { return (blockDim.x - 1) & ~(WARPSIZE - 1); }
+INLINE 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.
 INLINE int GetNumberOfWorkersInTeam() { return GetMasterThreadID(); }
 
 ////////////////////////////////////////////////////////////////////////////////
 // get thread id in team