[NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h

[JonChesterfield]

Summary

[NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h

Strictly there is one remaining difference wrt amdgcn - parallelLevel is
volatile qualified on amdgcn and not on nvptx. Determining whether this is
correct - and how to represent the different semantics of 'volatile' under
various conditions - is beyond the scope of this code motion patch.

[ABataev]

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

16

Do we need stdint.h include here?

34–37

Is this really target specific thing?

[JonChesterfield]

"Target specific" is suffering a little as an abstraction.

We have hardware specific (e.g. barrier implementation), language specific (e.g. attributes), driver specific (the above struct), vendor extensions (extra stuff one adds to the top level interface.h).

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

16

Sure. There's uint32_t and similar in this header, which could reasonably be the first or only one included in a translation unit.

34–37

This one is a driver/plugin specific thing. The hsa struct has another couple of fields.

[ABataev]

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

16

Then maybe include of stdint.h may be removed in omptarget-nvptx.h?

34–37

Why do you have different fields for hsa? Maybe we can use the same fields for NVPTX if they are useful?

[JonChesterfield]

• Remove stdint.h from omptarget-nvptx.h

[JonChesterfield]

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

34–37

The struct definition looks likely to be useful for nvptx, though the fields should probably then be wired up. Defined for gcn/hsa as:

struct omptarget_device_environmentTy {
  int32_t debug_level;   // gets value of envvar LIBOMPTARGET_DEVICE_RTL_DEBUG
                         // only useful for Debug build of deviceRTLs 
  int32_t num_devices;   // gets number of active offload devices 
  int32_t device_num;    // gets a value 0 to num_devices-1
};

where the fields are used by additional functions in libcall, e.g.

EXTERN int omp_get_num_devices(void) {
  PRINT(LD_IO, "call omp_get_num_devices() returns device_size %d\n",
        omptarget_device_environment.num_devices);
  return omptarget_device_environment.num_devices;
}

Would you like omp_get_num_devices / omp_get_device_num added to the generic interface? If so I'll write the patch, and reduce this diff to exclude the struct move.

[JonChesterfield]

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

16

Removed. Thanks

[JonChesterfield]

• Return device_environmentTy to original location, intending to implement num_devices calls for ptx

[ABataev]

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

34–37

Are they required by the standard?

[JonChesterfield]

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

34–37

I don't know. I'll find out before creating a diff which proposes said functions - in the meantime I've dropped it from this patch.

[ABataev]

LG

[JonChesterfield]

Committed, thanks.

Leaving this open to find out whether phabricator notices and auto-closes it.

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

 // and functions.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef __OMPTARGET_NVPTX_H
 #define __OMPTARGET_NVPTX_H
 
 // std includes
-#include <stdint.h>
-#include <stdlib.h>
-
 #include <inttypes.h>
-
-// cuda includes
-#include <cuda.h>
 #include <math.h>
+#include <stdlib.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
 
private:
   void **args;
   // starts off as MAX_SHARED_ARGS but can increase in size.
   uint32_t nArgs;
 };
 
 extern __device__ __shared__ omptarget_nvptx_SharedArgs
     omptarget_nvptx_globalArgs;
 
-// Data sharing related quantities, need to match what is used in the compiler.
-enum DATA_SHARING_SIZES {
-  // The maximum number of workers in a kernel.
-  DS_Max_Worker_Threads = 992,
-  // The size reserved for data in a shared memory slot.
-  DS_Slot_Size = 256,
-  // The slot size that should be reserved for a working warp.
-  DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
-  // The maximum number of warps in use
-  DS_Max_Warp_Number = 32,
-  // The size of the preallocated shared memory buffer per team
-  DS_Shared_Memory_Size = 128,
-};
-
 // Data structure to keep in shared memory that traces the current slot, stack,
 // and frame pointer as well as the active threads that didn't exit the current
 // environment.
 struct DataSharingStateTy {
   __kmpc_data_sharing_slot *SlotPtr[DS_Max_Warp_Number];
   void *StackPtr[DS_Max_Warp_Number];
   void * volatile FramePtr[DS_Max_Warp_Number];
   __kmpc_impl_lanemask_t ActiveThreads[DS_Max_Warp_Number];

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 <cuda.h>
 #include <stdint.h>
 
 #include "option.h"
 
+// Data sharing related quantities, need to match what is used in the compiler.
+enum DATA_SHARING_SIZES {
+  // The maximum number of workers in a kernel.
+  DS_Max_Worker_Threads = 992,
+  // The size reserved for data in a shared memory slot.
+  DS_Slot_Size = 256,
+  // The slot size that should be reserved for a working warp.
+  DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
+  // The maximum number of warps in use
+  DS_Max_Warp_Number = 32,
+  // The size of the preallocated shared memory buffer per team
+  DS_Shared_Memory_Size = 128,
+};
+
 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);