[libomptarget] Implement target_impl for amdgcn

[JonChesterfield]

Summary

[libomptarget] Implement target_impl for amdgcn

Smallest atomic addition for a new target. Implements enough of the amdgcn
specific code that some of the source files under nvptx/src could be compiled,
without modification, to run on amdgcn.

This foreshadows a work in progress patch to move said source out of nvptx/src.
Patch based on fork at https://github.com/ROCm-Developer-Tools/llvm-project

[ABataev]

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

24

__forceinline__?

82–83

Maybe ULL suffix instead of just L?

91

0xffffffffffffffffULL?

[JonChesterfield]

This is dead code, without any means of compiling it in tree. The clang patches aren't ready to go yet and will hopefully be replaced by the OpenMPIRBuilder.

I'd prefer to get the amdgcn deviceRTL in tree, preferably as incremental patches, so that I can use it while bringing up clang. It's not essential though. It's certainly reasonable to reject this patch as premature.

[JonChesterfield]

• Prefer uint64_c to trailing suffix

[JonChesterfield]

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

24

inline is sufficient when building with clang. It's possible the nvptx version should have an #ifdef that chooses between the two based on clang vs nvcc.

82–83

Nice spot. Don't really want the sign conversions here, fixed. Although to use UINT64_C instead of the suffix.

Strictly speaking that should probably be 32u (or UINT32_C(32), but that feels like a step too far) as well. Thoughts?

91

I'd prefer to use the UINT64_C notation. It's longer but it's harder to get wrong.

[ABataev]

LG

[JonChesterfield]

Great, thanks!

[jdoerfert]

I love it when a plan comes together

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

+//===--- amdgcn_interface.h - OpenMP interface definitions ------- 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _AMDGCN_INTERFACE_H_
+#define _AMDGCN_INTERFACE_H_
+
+#include <stdint.h>
+
+#define EXTERN extern "C" __attribute__((device))
+typedef uint64_t __kmpc_impl_lanemask_t;
+
+#endif

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

+//===------------ target_impl.h - AMDGCN 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_
+
+#ifndef __AMDGCN__
+#error "amdgcn target_impl.h expects to be compiled under __AMDGCN__"
+#endif
+
+#include <stdint.h>
+#include "amdgcn_interface.h"
+
+#define DEVICE __device__
+#define INLINE __inline__ DEVICE
+#define NOINLINE __noinline__ DEVICE
+
+////////////////////////////////////////////////////////////////////////////////
+// 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 64
+
+
+// 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.
+#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 maximum number of workers in a kernel.
+  DS_Max_Worker_Threads = 960,
+  // 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 = 16,
+};
+
+// warp vote function
+EXTERN uint64_t __ballot64(int predicate);
+// initialized with a 64-bit mask with bits set in positions less than the
+// thread's lane number in the warp
+EXTERN uint64_t __lanemask_lt();
+// initialized with a 64-bit mask with bits set in positions greater than the
+// thread's lane number in the warp
+EXTERN uint64_t __lanemask_gt();
+
+EXTERN void llvm_amdgcn_s_barrier();
+
+// CU id
+EXTERN unsigned __smid();
+
+INLINE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
+  lo = (uint32_t)(val & UINT64_C(0x00000000FFFFFFFF));
+  hi = (uint32_t)((val & UINT64_C(0xFFFFFFFF00000000)) >> 32);
+}
+
+INLINE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
+  return (((uint64_t)hi) << 32) | (uint64_t)lo;
+}
+
+static const __kmpc_impl_lanemask_t __kmpc_impl_all_lanes =
+    UINT64_C(0xffffffffffffffff);
+
+INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
+  return __lanemask_lt();
+}
+
+INLINE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
+  return __lanemask_gt();
+}
+
+INLINE uint32_t __kmpc_impl_smid() {
+  return __smid();
+}
+
+INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __ffsll(x); }
+
+INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __popcll(x); }
+
+INLINE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
+  return __ballot64(1);
+}
+
+INLINE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                     int32_t SrcLane) {
+  return __shfl(Var, SrcLane, WARPSIZE);
+}
+
+INLINE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
+                                          uint32_t Delta, int32_t Width) {
+  return __shfl_down(Var, Delta, Width);
+}
+
+INLINE void __kmpc_impl_syncthreads() { llvm_amdgcn_s_barrier(); }
+
+INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
+  // we have protected the master warp from releasing from its barrier
+  // due to a full workgroup barrier in the middle of a work function.
+  // So it is ok to issue a full workgroup barrier here.
+  __builtin_amdgcn_s_barrier();
+}
+
+#endif

openmp/libomptarget/deviceRTLs/interface.h

 //===----------------------------------------------------------------------===//
 
 #ifndef _INTERFACES_H_
 #define _INTERFACES_H_
 
 #include <stddef.h>
 #include <stdint.h>
 
+#ifdef __AMDGCN__
+#include "amdgcn/src/amdgcn_interface.h"
+#endif
 #ifdef __CUDACC__
 #include "nvptx/src/nvptx_interface.h"
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
 // OpenMP interface
 ////////////////////////////////////////////////////////////////////////////////