[libomptarget] Specialize amdgpu devicertl on wave size for gfx10

[JonChesterfield]

Summary

Use 32 bit arithmetic instead of relying on llvm to recognise
that the high half of various uint64_t values is zero for wave32 code.

Performance optimisation only. Relies on D108380 and D108391.

[JonChesterfield]

Low priority, posting it so I don't forget about it. Would remove the only reviewer but phab automatically re-adds you.

[JonChesterfield]

Only useful after D108708 has landed

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

144–149

This should be a 'runtime' switch on gridvalues, will update the patch

[JonChesterfield]

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

52

As a performance optimisation, this is probably in the noise.

However it will eliminate all the warp32 vs wave64 differences in the deviceRTL, making gfx10 a useful datapoint for debugging works on nvptx and fails on amdgpu. That is, if gfx10 works, it suggests the bug is in wave size. If it fails, it suggests the bug is not in wave size.

[JonChesterfield]

• less macros

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

 //===------- target_impl.h - AMDGCN OpenMP GPU implementation ----- HIP -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Declarations and definitions of target specific functions and constants
 //
 //===----------------------------------------------------------------------===//
 #ifndef OMPTARGET_AMDGCN_TARGET_IMPL_H
 #define OMPTARGET_AMDGCN_TARGET_IMPL_H
 
 #ifndef __AMDGCN__
 #error "amdgcn target_impl.h expects to be compiled under __AMDGCN__"
 #endif
 
 #include "amdgcn_interface.h"
 
 #include <stddef.h>
 #include <stdint.h>
 
 // subset of inttypes.h
 #define PRId64 "ld"
 #define PRIu64 "lu"
 
-typedef uint64_t __kmpc_impl_lanemask_t;
-
 #define INLINE inline
 #define NOINLINE __attribute__((noinline))
 #define ALIGN(N) __attribute__((aligned(N)))
 #define PLUGIN_ACCESSIBLE                                                      \
   __attribute__((used)) /* Don't discard values the plugin reads */            \
   __attribute__((visibility("default"))) /* Access via SHT_HASH */             \
   __attribute__((section(".data")))      /* Not .bss, can write before load */
 
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 INLINE constexpr const llvm::omp::GV &getGridValue() {
   return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
 }
 
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // The following def must match the absolute limit hardwired in the host RTL
 // max number of threads per team
 enum { MAX_THREADS_PER_TEAM = getGridValue().GV_Max_WG_Size };
 enum { WARPSIZE = getGridValue().GV_Warp_Size };
 
+namespace detail {
+template <unsigned> struct UnsignedToType;
+template <> struct UnsignedToType<64u> { using type = uint64_t; };
+template <> struct UnsignedToType<32u> { using type = uint32_t; };
+} // namespace detail
+
+using __kmpc_impl_lanemask_t = detail::UnsignedToType<WARPSIZE>::type;
+
 // Maximum number of omp state objects per SM allocated statically in global
 // memory.
 #define OMP_STATE_COUNT 32
 #define MAX_SM 64
 
 #define OMP_ACTIVE_PARALLEL_LEVEL 128
 
 // Data sharing related quantities, need to match what is used in the compiler.
 enum DATA_SHARING_SIZES {
   // The size reserved for data in a shared memory slot.
   DS_Slot_Size = getGridValue().GV_Slot_Size,
   // The slot size that should be reserved for a working warp.
   DS_Worker_Warp_Slot_Size = getGridValue().warpSlotSize(),
   // The maximum number of warps in use
   DS_Max_Warp_Number = getGridValue().maxWarpNumber(),
 };
 
 enum : __kmpc_impl_lanemask_t {
   __kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
 };
 
 // The return code of printf is not checked in the call sites in this library.
 // A call to a function named printf currently hits some special case handling
 // for opencl, which translates to calls that do not presently exist for openmp
 // Therefore, for now, stub out printf while building this library.
 #define printf(...)
 
 #endif

openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip

EXTERN double __kmpc_impl_get_wtime() {
   // This will probably need to be found by measurement on a number of
   // architectures. Until then, return 0, which is very inaccurate as a
   // timer but resolves the undefined symbol at link time.
   return 0;
 }
 
 // Warp vote function
 EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
+  static_assert(WARPSIZE == 64 || WARPSIZE == 32, "");
+  if (WARPSIZE == 64) {
     return __builtin_amdgcn_read_exec();
+  } else {
+    return __builtin_amdgcn_read_exec_lo();
+  }
 }
 
 static void pteam_mem_barrier(uint32_t num_threads, uint32_t *barrier_state) {
   __atomic_thread_fence(__ATOMIC_ACQUIRE);
 
   uint32_t num_waves = (num_threads + WARPSIZE - 1) / WARPSIZE;
 
   // Partial barrier implementation for amdgcn.
EXTERN int __kmpc_get_hardware_num_threads_in_block() {
   return get_workgroup_dim(__builtin_amdgcn_workgroup_id_x(),
                            __builtin_amdgcn_grid_size_x(),
                            __builtin_amdgcn_workgroup_size_x());
 }
 
 EXTERN unsigned GetWarpId() { return __kmpc_get_hardware_thread_id_in_block() / WARPSIZE; }
 EXTERN unsigned GetWarpSize() { return WARPSIZE; }
 EXTERN unsigned GetLaneId() {
+  static_assert(WARPSIZE == 64 || WARPSIZE == 32, "");
+  if (WARPSIZE == 64) {
     return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
+  } else {
+    return __builtin_amdgcn_mbcnt_lo(~0u, 0u);
+  }
 }
 
 EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads() {
   return __kmpc_get_hardware_num_threads_in_block();
 }
 
 // Atomics
 uint32_t __kmpc_atomic_add(uint32_t *Address, uint32_t Val) {