[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 ↗

(On Diff #323105)

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 ↗

(On Diff #323105)

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 ↗

(On Diff #323105)

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 ↗

(On Diff #323105)

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_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

 // }
 #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

 #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 "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

 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