[NFC] [libomptarget] Move option.h into target_impl.h, name a magic number

[JonChesterfield]

Summary

[NFC] [libomptarget] Move option.h into target_impl.h, name a magic number

Essentially everything under option.h is target dependent. This change moves the
existing code and names a mask that depends on the warp/wavefront size.

The macros could be changed to enums / functions / deleted subsequently.

[ABataev]

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

57–62

Is this a platform-dependent thing?

[JonChesterfield]

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

133

Note that this change from 0x1F to a function of warp size means the mask is 0x3F on 64 wide systems.

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

57–62

The inline/noinline macros are different for cuda and hip. Possibly also different for opencl, though I haven't checked.

True/false macros are things that happened to be in option.h. True is unused, false has a few uses related to debug printing. I'd like to remove those subsequently.

[ABataev]

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

133

Maybe it would be better to add 2 new functions, like getWarpId() and getThreadIdInWarp() or something like this and make them platform-dependent?

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

57–62

My question was not about inine/noinline, just about true/false stuff. They do not look like the platform dependent stuff. If you say that you can remove them, just remove them, no need to move it to the platform-specific header.

[JonChesterfield]

• Remove TRUE/FALSE macros
• add platform dependent functions to wrap warp size dependency

[JonChesterfield]

Not totally sure on __kmpcImplGetWarpId vs __kmpv_impl_get_warp_id, but suspect we should start to prefer the former.

[JonChesterfield]

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

133

Like it, thanks.

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

57–62

Also good, done.

[ABataev]

If everybody else is fine with the naming of the functions, the patch is good to go, I think.

[JonChesterfield]

In D69423#1721433, @ABataev wrote:

If everybody else is fine with the naming of the functions, the patch is good to go, I think.

Sticking with the local convention is probably less contentious, I think I'll swap it over. At some point we can run a find & replace across the rtl to change to the llvm style.

[JonChesterfield]

• use local convention for naming, clang-format changed functions

[ABataev]

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

135–141

Maybe, split this patch into 2 and introduce these functions and new processing of check functions in the new patch?

[JonChesterfield]

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

135–141

I'm not sure we gain much by dividing this patch. If you like I could split it into something like:

• delete macros
• move constants into target_impl
• introduce & use one function
• introduce & use the other function

I don't have time to do that just now but will be able to later

[ABataev]

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

135–141

2 last can be merged into 1. All other steps look good.

[JonChesterfield]

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

135–141

Cool. Will do so later today / early tomorrow. Cheers

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

 //===----------------------------------------------------------------------===//
 
 #include "omptarget-nvptx.h"
 
 EXTERN int32_t __kmpc_cancellationpoint(kmp_Ident *loc, int32_t global_tid,
                                         int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancellationpoint(cancel val %d)\n", (int)cancelVal);
   // disabled
-  return FALSE;
+  return 0;
 }
 
 EXTERN int32_t __kmpc_cancel(kmp_Ident *loc, int32_t global_tid,
                              int32_t cancelVal) {
   PRINT(LD_IO, "call kmpc_cancel(cancel val %d)\n", (int)cancelVal);
   // disabled
-  return FALSE;
+  return 0;
 }

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 "option.h"
+#include "target_impl.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)(__kmpc_impl_get_warp_id()),
+         (int)(__kmpc_impl_get_thread_id_in_warp()), 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)(__kmpc_impl_get_warp_id()),
+           (int)(__kmpc_impl_get_thread_id_in_warp()), parameters...);
   assert(cond);
 }
 
 NOINLINE static void check(bool cond) { assert(cond); }
 #endif
 
 // set flags that are tested (inclusion properties)
 

 #define PRINT(_flag, _str, _args...)                                           \
   {                                                                            \
     if (omptarget_device_environment.debug_level && DON(_flag)) {              \
       log("<b %2d, t %4d, w %2d, l %2d>: " _str, _args);                       \
     }                                                                          \
   }
 #else
 
-#define DON(_flag) (FALSE)
+#define DON(_flag) (0)
 #define PRINT0(flag, str)
 #define PRINT(flag, str, _args...)
 
 #endif
 
 // for printing without worring about precision, pointers...
 #define P64(_x) ((unsigned long long)(_x))
 
#define ASSERT(_flag, _cond, _str, _args...)                                   \
     if (TON(_flag)) {                                                          \
       check((_cond), "<b %3d, t %4d, w %2d, l %d2> ASSERT: " _str "\n",        \
             _args);                                                            \
     }                                                                          \
   }
 
 #else
 
-#define TON(_flag) (FALSE)
+#define TON(_flag) (0)
 #define ASSERT0(_flag, _cond, _str)
 #define ASSERT(_flag, _cond, _str, _args...)
 
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // early defs for warning
 
#define WARNING(_flag, _str, _args...)                                         \
   {                                                                            \
     if (WON(_flag)) {                                                          \
       log("<b %2d, t %4d, w %2d, l %2d> WARNING: " _str, _args);               \
     }                                                                          \
   }
 
 #else
 
-#define WON(_flag) (FALSE)
+#define WON(_flag) (0)
 #define WARNING0(_flag, _str)
 #define WARNING(_flag, _str, _args...)
 
 #endif
 
 #endif

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

EXTERN int omp_get_partition_num_places(void) {
   return 0;
 }
 
 EXTERN void omp_get_partition_place_nums(int *place_nums) {
   PRINT0(LD_IO, "call to omp_get_partition_place_nums()\n");
 }
 
 EXTERN int omp_get_cancellation(void) {
-  int rc = FALSE; // currently false only
+  int rc = 0;
   PRINT(LD_IO, "call omp_get_cancellation() returns %d\n", rc);
   return rc;
 }
 
 EXTERN void omp_set_default_device(int deviceId) {
   PRINT0(LD_IO, "call omp_get_default_device() is undef on device\n");
 }
 

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

INLINE static void for_static_init(int32_t gtid, int32_t schedtype,
     }
     case kmp_sched_distr_static_chunk_sched_static_chunkone: {
       ForStaticChunk(lastiter, lb, ub, stride, chunk,
                      numberOfActiveOMPThreads * GetOmpTeamId() + gtid,
                      GetNumberOfOmpTeams() * numberOfActiveOMPThreads);
       break;
     }
     default: {
-      ASSERT(LT_FUSSY, FALSE, "unknown schedtype %d", (int)schedtype);
+      ASSERT(LT_FUSSY, 0, "unknown schedtype %d", (int)schedtype);
       PRINT(LD_LOOP, "unknown schedtype %d, revert back to static chunk\n",
             (int)schedtype);
       ForStaticChunk(lastiter, lb, ub, stride, chunk, gtid,
                      numberOfActiveOMPThreads);
       break;
     }
     }
     // copy back

openmp/libomptarget/deviceRTLs/nvptx/src/omptarget-nvptx.h

 
 #include <inttypes.h>
 
 // cuda includes
 #include <cuda.h>
 #include <math.h>
 
 // local includes
+#include "target_impl.h"
 #include "debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
-#include "option.h"    // choices we have
 #include "state-queue.h"
 #include "support.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1
 
 // used by the library for the interface with the app
 #define DISPATCH_FINISHED 0
 #define DISPATCH_NOTFINISHED 1

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

-//===------------ option.h - NVPTX OpenMP GPU options ------------ 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
-//
-//===----------------------------------------------------------------------===//
-//
-// GPU default options
-//
-//===----------------------------------------------------------------------===//
-#ifndef _OPTION_H_
-#define _OPTION_H_
-
-#include "interface.h"
-
-////////////////////////////////////////////////////////////////////////////////
-// Kernel options
-////////////////////////////////////////////////////////////////////////////////
-
-////////////////////////////////////////////////////////////////////////////////
-// The following def must match the absolute limit hardwired in the host RTL
-// max number of threads per team
-#define MAX_THREADS_PER_TEAM 1024
-
-#define WARPSIZE 32
-
-// The named barrier for active parallel threads of a team in an L1 parallel
-// region to synchronize with each other.
-#define L1_BARRIER (1)
-
-// Maximum number of preallocated arguments to an outlined parallel/simd function.
-// Anything more requires dynamic memory allocation.
-#define MAX_SHARED_ARGS 20
-
-// Maximum number of omp state objects per SM allocated statically in global
-// memory.
-#if __CUDA_ARCH__ >= 700
-#define OMP_STATE_COUNT 32
-#define MAX_SM 84
-#elif __CUDA_ARCH__ >= 600
-#define OMP_STATE_COUNT 32
-#define MAX_SM 56
-#else
-#define OMP_STATE_COUNT 16
-#define MAX_SM 16
-#endif
-
-#define OMP_ACTIVE_PARALLEL_LEVEL 128
-
-////////////////////////////////////////////////////////////////////////////////
-// algo options
-////////////////////////////////////////////////////////////////////////////////
-
-////////////////////////////////////////////////////////////////////////////////
-// misc options (by def everythig here is device)
-////////////////////////////////////////////////////////////////////////////////
-
-#define INLINE __forceinline__ __device__
-#define NOINLINE __noinline__ __device__
-#ifndef TRUE
-#define TRUE 1
-#endif
-#ifndef FALSE
-#define FALSE 0
-#endif
-
-#endif

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

 }
 
 // Do nothing. The host guarantees we started the requested number of
 // teams and we only need inspection of gridDim.
 
 EXTERN void __kmpc_push_num_teams(kmp_Ident *loc, int32_t tid,
                                   int32_t num_teams, int32_t thread_limit) {
   PRINT(LD_IO, "call kmpc_push_num_teams %d\n", (int)num_teams);
-  ASSERT0(LT_FUSSY, FALSE,
+  ASSERT0(LT_FUSSY, 0,
           "should never have anything with new teams on device");
 }
 
 EXTERN void __kmpc_push_proc_bind(kmp_Ident *loc, uint32_t tid,
                                   int proc_bind) {
   PRINT(LD_IO, "call kmpc_push_proc_bind %d\n", (int)proc_bind);
 }

openmp/libomptarget/deviceRTLs/nvptx/src/state-queue.h

 //
 //===----------------------------------------------------------------------===//
 
 #ifndef __STATE_QUEUE_H
 #define __STATE_QUEUE_H
 
 #include <stdint.h>
 
-#include "option.h" // choices we have
+#include "target_impl.h"
 
 template <typename ElementType, uint32_t SIZE> class omptarget_nvptx_Queue {
 private:
   ElementType elements[SIZE];
   volatile ElementType *elementQueue[SIZE];
   volatile uint32_t head;
   volatile uint32_t ids[SIZE];
   volatile uint32_t tail;

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

 //===------------ target_impl.h - NVPTX OpenMP GPU options ------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 #ifndef _TARGET_IMPL_H_
 #define _TARGET_IMPL_H_
 
 #include <stdint.h>
 
-#include "option.h"
+////////////////////////////////////////////////////////////////////////////////
+// Kernel options
+////////////////////////////////////////////////////////////////////////////////
+
+////////////////////////////////////////////////////////////////////////////////
+// The following def must match the absolute limit hardwired in the host RTL
+// max number of threads per team
+#define MAX_THREADS_PER_TEAM 1024
+
+#define WARPSIZE 32
+
+// The named barrier for active parallel threads of a team in an L1 parallel
+// region to synchronize with each other.
+#define L1_BARRIER (1)
+
+// Maximum number of preallocated arguments to an outlined parallel/simd function.
+// Anything more requires dynamic memory allocation.
+#define MAX_SHARED_ARGS 20
+
+// Maximum number of omp state objects per SM allocated statically in global
+// memory.
+#if __CUDA_ARCH__ >= 700
+#define OMP_STATE_COUNT 32
+#define MAX_SM 84
+#elif __CUDA_ARCH__ >= 600
+#define OMP_STATE_COUNT 32
+#define MAX_SM 56
+#else
+#define OMP_STATE_COUNT 16
+#define MAX_SM 16
+#endif
+
+#define OMP_ACTIVE_PARALLEL_LEVEL 128
+
+#define INLINE __forceinline__ __device__
+#define NOINLINE __noinline__ __device__
 
 INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "l"(val));
 }
 
 INLINE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   uint64_t val;
   asm volatile("mov.b64 %0, {%1,%2};" : "=l"(val) : "r"(lo), "r"(hi));
   return val;
 }
 
 static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
     UINT32_C(0xffffffff);
 
 INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
   __kmpc_impl_lanemask_t res;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(res));
   return res;

 INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
 #if CUDA_VERSION >= 9000
   __syncwarp(Mask);
 #else
   // In Cuda < 9.0 no need to sync threads in warps.
 #endif // CUDA_VERSION
 }
 
+INLINE uint32_t __kmpc_impl_get_warp_id() { return threadIdx.x / WARPSIZE; }
+
+INLINE uint32_t __kmpc_impl_get_thread_id_in_warp() {
+  int threadIdxMask = WARPSIZE - 1;
+  return threadIdx.x & threadIdxMask;
+}
+
 #endif