[libomptarget][amdgcn] Build amdgcn devicertl as openmp

[JonChesterfield]

Summary

[libomptarget][amdgcn] Build amdgcn devicertl as openmp

Change cmake to build as openmp and fix up some minor errors in the code.

[JonChesterfield]

This doesn't toggle the cmake to build the amdgcn plugin by default. I'm planning to do that after the dust settles in the internal builds from this.

[tianshilei1992]

In D96533#2557822, @JonChesterfield wrote:

This doesn't toggle the cmake to build the amdgcn plugin by default. I'm planning to do that after the dust settles in the internal builds from this.

One nice thing if AMD can build deviceRTLs with OpenMP is, we can completely get rid of the macro DEVICE.

[JonChesterfield]

Yep. I'm walking a slightly fine line between this and the internal rocm builds so would like to keep the functional changes separate from the cleanup. Plan is ship this, which 'should' work internally, pick up any pieces if it doesn't, then drop the DEVICE and probably INLINE macros. Maybe get rid of some of EXTERN as well.

Curiously the internal build does compile this library as openmp, and has done since December, but uses a slightly different clang and a slightly different devicertl.

[jdoerfert]

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

46

Do we need to declare builtins? I guess we can but we should not have to, right?

[JonChesterfield]

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

46

Today, yes. I think the 'right' behaviour is to allow these within an appropriate declare variant, and an adequate approximation would be to allow them everywhere.

As raw openmp, the assumption is that any function can be called from the host. That doesn't mean anything for __builtin_amdgcn_dispatch_ptr, so it's reasonable that these aren't automatically available.

At present though, I'm just trying to get this library to compile with trunk clang as-is. Curiously, if one of these is declared with the wrong type, clang emits an error on the declaration.

[jdoerfert]

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

46

I can just call these like other builtins (when I target amdgcn):
https://godbolt.org/z/nabGT7

[JonChesterfield]

• drop unused declaration

[JonChesterfield]

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

46

That is interesting.

I wrote this patch against a version of --cuda-device-only. These may have only raised an error on the host pass, which is not invoked by the present -Xclang themed cmake.

[JonChesterfield]

• Drop amdgcn intrinsics, not needed with host pass disabled

[jdoerfert]

LGTM

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

 # for both in-tree and out-of-tree build
 if (NOT CMAKE_ARCHIVE_OUTPUT_DIRECTORY)
   set(OUTPUTDIR ${CMAKE_CURRENT_BINARY_DIR})
 else()
   set(OUTPUTDIR ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
 endif()
 
 # create libraries
-set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900)
+set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900 gfx906)
 if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
   set(mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
 endif()
 
 macro(add_cuda_bc_library)
   set(cu_cmd ${AOMP_BINDIR}/clang++
+    -xc++
+    -c
     -std=c++14
-    -fcuda-rdc
+    -target amdgcn
+    -emit-llvm
+    -Xclang -aux-triple -Xclang x86_64-unknown-linux-gnu # see nvptx
+    -fopenmp -fopenmp-cuda-mode -Xclang -fopenmp-is-device
+    -D__AMDGCN__
+    -Xclang -target-cpu -Xclang ${mcpu}
     -fvisibility=default
-    --cuda-device-only
     -Wno-unused-value
-    -x hip
-    -nogpulib -nogpuinc
+    -nogpulib
     -O${optimization_level}
-    --cuda-gpu-arch=${mcpu}
     ${CUDA_DEBUG}
     -I${CMAKE_CURRENT_SOURCE_DIR}/src
     -I${devicertl_base_directory})
 
   set(bc1_files)
 
   foreach(file ${ARGN})
     get_filename_component(fname ${file} NAME_WE)

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))
+#define EXTERN extern "C"
 typedef uint64_t __kmpc_impl_lanemask_t;
 typedef uint32_t omp_lock_t; /* arbitrary type of the right length */
 
 EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads();
 
 #endif

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

+//===--- amdgcn_intrinsics.h - Intrinsics used by deviceRTL ---------------===//
+//
+// 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_INTRINSICS_H_
+#define _AMDGCN_INTRINSICS_H_
+
+#include <stdint.h>
+
+#include "target_impl.h"
+
+#ifdef _OPENMP
+// Openmp doesn't pull these builtins into scope, but does error if the type is
+// incorrect
+// This may be a quirk of openmp's compile for host + device assumption, where
+// these don't resolve to anything on the host
+
+EXTERN uint32_t __builtin_amdgcn_atomic_inc32(volatile uint32_t *, uint32_t,
+                                              uint32_t, const char *);
+EXTERN void __builtin_amdgcn_s_barrier(void);
+EXTERN void __builtin_amdgcn_fence(uint32_t, const char *);
+
+EXTERN void __builtin_amdgcn_s_sleep(int);
+
+EXTERN uint32_t __builtin_amdgcn_workitem_id_x(void);
+EXTERN uint32_t __builtin_amdgcn_workgroup_id_x(void);
+EXTERN uint16_t __builtin_amdgcn_workgroup_size_x(void);
+EXTERN uint32_t __builtin_amdgcn_grid_size_x(void);
+
+EXTERN uint64_t __builtin_amdgcn_s_memrealtime(void);
+EXTERN uint32_t __builtin_amdgcn_s_getreg(int32_t);
+EXTERN uint64_t __builtin_amdgcn_read_exec(void);
+
+EXTERN __attribute__((address_space(4))) void *
+__builtin_amdgcn_dispatch_ptr() noexcept;
+
+EXTERN uint32_t __builtin_amdgcn_mbcnt_lo(uint32_t, uint32_t);
+EXTERN uint32_t __builtin_amdgcn_mbcnt_hi(uint32_t, uint32_t);
+EXTERN int32_t __builtin_amdgcn_ds_bpermute(int32_t, int32_t);
+#endif
+
+#endif

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

 static DEVICE void warn() {
   PRINT0(LD_ALL, "Locks are not supported in this thread mapping model");
 }
 
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_set_lock(omp_lock_t *) { warn(); }
 DEVICE void __kmpc_impl_unset_lock(omp_lock_t *) { warn(); }
-DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock) { warn(); }
+DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock) { warn(); return 0;}
 
 #pragma omp end declare target

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

 //===-------- amdgcn_smid.hip - AMDGCN smid 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
 //
 //===----------------------------------------------------------------------===//
 #pragma omp declare target
 
 #include "target_impl.h"
+#include "amdgcn_intrinsics.h"
 
 // Partially derived fom hcc_detail/device_functions.h
 
 // HW_ID Register bit structure
 // WAVE_ID     3:0     Wave buffer slot number. 0-9.
 // SIMD_ID     5:4     SIMD which the wave is assigned to within the CU.
 // PIPE_ID     7:6     Pipeline from which the wave was dispatched.
 // CU_ID       11:8    Compute Unit the wave is assigned to.

 // }
 #define ENCODE_HWREG(WIDTH, OFF, REG) (REG | (OFF << 6) | ((WIDTH - 1) << 11))
 
 // Note: The results can be changed by a context switch
 // Return value in [0 2^SE_ID_SIZE * 2^CU_ID_SIZE), which is an upper
 // bound on how many compute units are available. Some values in this
 // range may never be returned if there are fewer than 2^CU_ID_SIZE CUs.
 
-DEVICE uint32_t __kmpc_impl_smid() {
+EXTERN uint32_t __kmpc_impl_smid() {
   uint32_t cu_id = __builtin_amdgcn_s_getreg(
       ENCODE_HWREG(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID));
   uint32_t se_id = __builtin_amdgcn_s_getreg(
       ENCODE_HWREG(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID));
   return (se_id << HW_ID_CU_ID_SIZE) + cu_id;
 }
 
 #pragma omp end declare target

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 <inttypes.h>
 #include <stddef.h>
 #include <stdint.h>
 
-#define DEVICE __attribute__((device))
+#define DEVICE
 #define INLINE inline DEVICE
 #define NOINLINE __attribute__((noinline)) DEVICE
-#define SHARED(NAME) __attribute__((shared)) NAME
-#define EXTERN_SHARED(NAME) __attribute__((shared)) NAME
 #define ALIGN(N) __attribute__((aligned(N)))
 
 ////////////////////////////////////////////////////////////////////////////////
 // Kernel options
 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 // The following def must match the absolute limit hardwired in the host RTL

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

 //===------- target_impl.hip - 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Definitions of target specific functions
 //
 //===----------------------------------------------------------------------===//
 #pragma omp declare target
 
+#include "amdgcn_intrinsics.h"
+#include "common/omptarget.h"
 #include "target_impl.h"
+#include "target_interface.h"
 
 // Implementations initially derived from hcc
 
 // Initialized with a 64-bit mask with bits set in positions less than the
 // thread's lane number in the warp
 EXTERN __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
   uint32_t lane = GetLaneId();
   int64_t ballot = __kmpc_impl_activemask();

 EXTERN int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t var,
                                           uint32_t laneDelta, int32_t width) {
   int self = GetLaneId();
   int index = self + laneDelta;
   index = (int)(laneDelta + (self & (width - 1))) >= width ? self : index;
   return __builtin_amdgcn_ds_bpermute(index << 2, var);
 }
 
-static DEVICE SHARED uint32_t L1_Barrier;
+uint32_t __kmpc_L1_Barrier [[clang::loader_uninitialized]];
+#pragma allocate(__kmpc_L1_Barrier) allocator(omp_pteam_mem_alloc)
 
 EXTERN void __kmpc_impl_target_init() {
   // Don't have global ctors, and shared memory is not zero init
-  __atomic_store_n(&L1_Barrier, 0u, __ATOMIC_RELEASE);
+  __atomic_store_n(&__kmpc_L1_Barrier, 0u, __ATOMIC_RELEASE);
 }
 
 EXTERN void __kmpc_impl_named_sync(uint32_t num_threads) {
   __atomic_thread_fence(__ATOMIC_ACQUIRE);
 
   uint32_t num_waves = num_threads / WARPSIZE;
 
   // Partial barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
   // reached the barrier, and one to count how many times the barrier has been
   // passed. These are packed in a single atomically accessed 32 bit integer.
   // Low bits for the number of waves, assumed zero before this call.
   // High bits to count the number of times the barrier has been passed.
 
   // precondition: num_waves != 0;
   // invariant: num_waves * WARPSIZE == num_threads;
   // precondition: num_waves < 0xffffu;
 
   // Increment the low 16 bits once, using the lowest active thread.
   uint64_t lowestActiveThread = __kmpc_impl_ffs(__kmpc_impl_activemask()) - 1;
   bool isLowest = GetLaneId() == lowestActiveThread;
 
   if (isLowest) {
-    uint32_t load =
-        __atomic_fetch_add(&L1_Barrier, 1, __ATOMIC_RELAXED); // commutative
+    uint32_t load = __atomic_fetch_add(&__kmpc_L1_Barrier, 1,
+                                       __ATOMIC_RELAXED); // commutative
 
     // Record the number of times the barrier has been passed
     uint32_t generation = load & 0xffff0000u;
 
     if ((load & 0x0000ffffu) == (num_waves - 1)) {
       // Reached num_waves in low bits so this is the last wave.
       // Set low bits to zero and increment high bits
       load += 0x00010000u; // wrap is safe
       load &= 0xffff0000u; // because bits zeroed second
 
       // Reset the wave counter and release the waiting waves
-      __atomic_store_n(&L1_Barrier, load, __ATOMIC_RELAXED);
+      __atomic_store_n(&__kmpc_L1_Barrier, load, __ATOMIC_RELAXED);
     } else {
       // more waves still to go, spin until generation counter changes
       do {
         __builtin_amdgcn_s_sleep(0);
-        load = __atomic_load_n(&L1_Barrier, __ATOMIC_RELAXED);
+        load = __atomic_load_n(&__kmpc_L1_Barrier, __ATOMIC_RELAXED);
       } while ((load & 0xffff0000u) == generation);
     }
   }
   __atomic_thread_fence(__ATOMIC_RELEASE);
 }
 
 namespace {
 DEVICE uint32_t get_grid_dim(uint32_t n, uint16_t d) {

openmp/libomptarget/deviceRTLs/common/omptarget.h

 //===---- omptarget.h - OpenMP GPU initialization ---------------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file contains the declarations of all library macros, types,
 // and functions.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_H
 #define OMPTARGET_H
 
 #include "common/allocator.h"
 #include "common/debug.h" // debug
 #include "common/state-queue.h"
 #include "common/support.h"
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "target_impl.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1
 

 getMyTopTaskDescriptor(bool isSPMDExecutionMode);
 INLINE omptarget_nvptx_TaskDescr *getMyTopTaskDescriptor(int globalThreadId);
 
 ////////////////////////////////////////////////////////////////////////////////
 // inlined implementation
 ////////////////////////////////////////////////////////////////////////////////
 
 INLINE uint32_t __kmpc_impl_ffs(uint32_t x) { return __builtin_ffs(x); }
-
 INLINE uint32_t __kmpc_impl_popc(uint32_t x) { return __builtin_popcount(x); }
+INLINE uint32_t __kmpc_impl_ffs(uint64_t x) { return __builtin_ffsl(x); }
+INLINE uint32_t __kmpc_impl_popc(uint64_t x) { return __builtin_popcountl(x); }
 
 #include "common/omptargeti.h"
 
 #endif