[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 ↗

(On Diff #226575)

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

[ABataev]

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

126 ↗

(On Diff #226575)

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

[JonChesterfield]

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

126 ↗

(On Diff #226575)

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/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)blockIdx.x, (int)threadIdx.x,
+         (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...);
+           (int)(GetWarpId()),
+           (int)(GetLaneId()), parameters...);
   assert(cond);
 }
 
 NOINLINE static void check(bool cond) { assert(cond); }
 #endif
 
 // set flags that are tested (inclusion properties)