[libomptarget][amdgcn] Build amdgcn devicertl as openmp
[libomptarget][amdgcn] Build amdgcn devicertl as openmp
Change cmake to build as openmp and fix up some minor errors in the code.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D96533
GitOrigin-RevId: 6f04addc8b2eee0d88b97facfa4fb7424b4b74bd
diff --git a/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt b/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
index 8d9abe5..0e890b1 100644
--- a/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
+++ b/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
@@ -90,22 +90,26 @@
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})
diff --git a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_interface.h b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_interface.h
index 80409d6..1e653da 100644
--- a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_interface.h
+++ b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_interface.h
@@ -11,7 +11,7 @@
#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 */
diff --git a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip
index f537fb2..77746fa 100644
--- a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip
+++ b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_locks.hip
@@ -26,6 +26,6 @@
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
diff --git a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip
index c850455..f18f8b5 100644
--- a/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip
+++ b/libomptarget/deviceRTLs/amdgcn/src/amdgcn_smid.hip
@@ -53,7 +53,7 @@
// 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(
diff --git a/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
index fb41332..6a40c18 100644
--- a/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ b/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
@@ -22,11 +22,9 @@
#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)))
////////////////////////////////////////////////////////////////////////////////
diff --git a/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip b/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
index cf04b48..51eaf71 100644
--- a/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
+++ b/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
@@ -11,7 +11,9 @@
//===----------------------------------------------------------------------===//
#pragma omp declare target
+#include "common/omptarget.h"
#include "target_impl.h"
+#include "target_interface.h"
// Implementations initially derived from hcc
@@ -66,11 +68,12 @@
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) {
@@ -94,8 +97,8 @@
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;
@@ -107,12 +110,12 @@
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);
}
}
diff --git a/libomptarget/deviceRTLs/common/omptarget.h b/libomptarget/deviceRTLs/common/omptarget.h
index 7692233..e179ca5 100644
--- a/libomptarget/deviceRTLs/common/omptarget.h
+++ b/libomptarget/deviceRTLs/common/omptarget.h
@@ -337,8 +337,9 @@
////////////////////////////////////////////////////////////////////////////////
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"
