[OpenMP][deviceRTLs] Separate declaration of target dependent functions from `target_impl.h`

[tianshilei1992]

Summary

This patch created a new header file target_interface.h for declarations of all target dependent functions. All future targets can get things work by simply implementing all functions declared in the header and macros/data same as each target_impl.h.

[JonChesterfield]

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

69–70

All the declarations that moved to targwt_interface should be deleted from here (as well as from nvptx target_impl)

openmp/libomptarget/deviceRTLs/common/omptarget.h

333

Just INLINE?

openmp/libomptarget/deviceRTLs/target_interface.h

28

I guess these should be EXTERN, not extern DEVICE

[JonChesterfield]

'new header file target_impl.h' should be 'new header file target_interface.h' in the commit message

[tianshilei1992]

openmp/libomptarget/deviceRTLs/common/omptarget.h

333

otherwise?

openmp/libomptarget/deviceRTLs/target_interface.h

28

They're C++ functions, so cannot use EXTERN directly which is expanded to extern "C" DEVICE.

[tianshilei1992]

Rebased and fixed comments

[tianshilei1992]

rebased

[JonChesterfield]

Couple of minor points about redundant function annotations, otherwise good. Thanks for moving the amdgpu code out of the header.

I wonder if the kmpc_impl functions should be consistently extern C, at least when not overloaded.

openmp/libomptarget/deviceRTLs/common/omptarget.h

333

I mean, INLINE expands to a definition that includes DEVICE, so we don't need both

openmp/libomptarget/deviceRTLs/target_interface.h

28

Ah, so 'extern' here is dead. Functions are non-static by default.

[tianshilei1992]

Updated comments

[tianshilei1992]

rebase to triger a new round of test

[tianshilei1992]

rebased

[JonChesterfield]

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

192

I missed the language linkage change from c++ to c on these, fixed on main.

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

enum DATA_SHARING_SIZES {
   DS_Worker_Warp_Slot_Size = WARPSIZE * DS_Slot_Size,
   // The maximum number of warps in use
   DS_Max_Warp_Number = 16,
 };
 
 enum : __kmpc_impl_lanemask_t {
   __kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
 };
 
-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;
-}
-
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
-DEVICE uint32_t __kmpc_impl_smid();
-DEVICE double __kmpc_impl_get_wtick();
-DEVICE double __kmpc_impl_get_wtime();
-
-INLINE uint64_t __kmpc_impl_ffs(uint64_t x) { return __builtin_ffsl(x); }
-INLINE uint64_t __kmpc_impl_popc(uint64_t x) { return __builtin_popcountl(x); }
-
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask();
-
-DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                     int32_t SrcLane);
-
-DEVICE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t, int32_t Var,
-                                          uint32_t Delta, int32_t Width);
-
-INLINE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
-
-INLINE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) {
-  // AMDGCN doesn't need to sync threads in a warp
-}
-
-// AMDGCN specific kernel initialization
-DEVICE void __kmpc_impl_target_init();
-
-// Equivalent to ptx bar.sync 1. Barrier until num_threads arrive.
-DEVICE void __kmpc_impl_named_sync(uint32_t num_threads);
-
-INLINE void __kmpc_impl_threadfence() {
-  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
-}
-
-INLINE void __kmpc_impl_threadfence_block() {
-  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
-}
-
-INLINE void __kmpc_impl_threadfence_system() {
-  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
-}
-
-// Calls to the AMDGCN layer (assuming 1D layout)
-INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
-INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
-DEVICE int GetNumberOfBlocksInKernel();
-DEVICE int GetNumberOfThreadsInBlock();
-DEVICE unsigned GetWarpId();
-DEVICE unsigned GetLaneId();
-
-// Atomics
-DEVICE uint32_t __kmpc_atomic_add(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_inc(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_max(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_exchange(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_cas(uint32_t *, uint32_t, uint32_t);
-
-static_assert(sizeof(unsigned long long) == sizeof(uint64_t), "");
-DEVICE unsigned long long __kmpc_atomic_exchange(unsigned long long *,
-                                                 unsigned long long);
-DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *,
-                                            unsigned long long);
-
-// Locks
-DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_set_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_unset_lock(omp_lock_t *lock);
-DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock);
-
-// Memory
-DEVICE void *__kmpc_impl_malloc(size_t x);
-DEVICE void __kmpc_impl_free(void *x);
-
 // DEVICE versions of part of libc
 INLINE void __assert_fail(const char *, const char *, unsigned int,
                           const char *) {
   __builtin_trap();
 }
 EXTERN int printf(const char *, ...);
 
 #endif

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

DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *Address,
                                             unsigned long long Val) {
   return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 
 // Stub implementations
 DEVICE void *__kmpc_impl_malloc(size_t) { return nullptr; }
 DEVICE void __kmpc_impl_free(void *) {}
 
+DEVICE 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);
+}
+
+DEVICE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
+  return (((uint64_t)hi) << 32) | (uint64_t)lo;
+}
+
+DEVICE void __kmpc_impl_syncthreads() { __builtin_amdgcn_s_barrier(); }
+
+DEVICE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t) {
+  // AMDGCN doesn't need to sync threads in a warp
+}
+
+DEVICE void __kmpc_impl_threadfence() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+}
+
+DEVICE void __kmpc_impl_threadfence_block() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+}
+
+DEVICE void __kmpc_impl_threadfence_system() {
+  __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+}
+
+// Calls to the AMDGCN layer (assuming 1D layout)
+DEVICE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
+DEVICE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
+
 #pragma omp end declare target

openmp/libomptarget/deviceRTLs/common/debug.h

 //   WARNING(flag, str, args)      : if warning flag is on, print the warning
 //   WON(flag)                     : return true if warning flag is on
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef _OMPTARGET_NVPTX_DEBUG_H_
 #define _OMPTARGET_NVPTX_DEBUG_H_
 
 #include "common/device_environment.h"
+#include "target_interface.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // set desired level of debugging
 ////////////////////////////////////////////////////////////////////////////////
 
 #define LD_SET_NONE 0ULL /* none */
 #define LD_SET_ALL -1ULL /* all */
 

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 "target_impl.h"
 #include "common/debug.h"     // debug
 #include "interface.h" // interfaces with omp, compiler, and user
 #include "common/state-queue.h"
 #include "common/support.h"
 
 #define OMPTARGET_NVPTX_VERSION 1.1
 
 // used by the library for the interface with the app

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

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

enum DATA_SHARING_SIZES {
   // The size of the preallocated shared memory buffer per team
   DS_Shared_Memory_Size = 128,
 };
 
 enum : __kmpc_impl_lanemask_t {
   __kmpc_impl_all_lanes = ~(__kmpc_impl_lanemask_t)0
 };
 
-DEVICE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi);
-DEVICE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi);
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
-DEVICE uint32_t __kmpc_impl_smid();
-DEVICE double __kmpc_impl_get_wtick();
-DEVICE double __kmpc_impl_get_wtime();
-
-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); }
-
 #ifndef CUDA_VERSION
 #error CUDA_VERSION macro is undefined, something wrong with cuda.
 #endif
 
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask();
-
-DEVICE int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,
-                                     int32_t SrcLane);
-
-DEVICE int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t Mask,
-                                          int32_t Var, uint32_t Delta,
-                                          int32_t Width);
-
-DEVICE void __kmpc_impl_syncthreads();
-DEVICE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask);
-
-// NVPTX specific kernel initialization
-DEVICE void __kmpc_impl_target_init();
-
-// Barrier until num_threads arrive.
-DEVICE void __kmpc_impl_named_sync(uint32_t num_threads);
-
-DEVICE void __kmpc_impl_threadfence();
-DEVICE void __kmpc_impl_threadfence_block();
-DEVICE void __kmpc_impl_threadfence_system();
-
-// Calls to the NVPTX layer (assuming 1D layout)
-DEVICE int GetThreadIdInBlock();
-DEVICE int GetBlockIdInKernel();
-DEVICE int GetNumberOfBlocksInKernel();
-DEVICE int GetNumberOfThreadsInBlock();
-DEVICE unsigned GetWarpId();
-DEVICE unsigned GetLaneId();
-
-// Atomics
-DEVICE uint32_t __kmpc_atomic_add(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_inc(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_max(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_exchange(uint32_t *, uint32_t);
-DEVICE uint32_t __kmpc_atomic_cas(uint32_t *, uint32_t, uint32_t);
-
-static_assert(sizeof(unsigned long long) == sizeof(uint64_t), "");
-DEVICE unsigned long long __kmpc_atomic_exchange(unsigned long long *,
-                                                 unsigned long long);
-DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *,
-                                            unsigned long long);
-
-// Locks
-DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_set_lock(omp_lock_t *lock);
-DEVICE void __kmpc_impl_unset_lock(omp_lock_t *lock);
-DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock);
-
-// Memory
-DEVICE void *__kmpc_impl_malloc(size_t);
-DEVICE void __kmpc_impl_free(void *);
-
 #endif

openmp/libomptarget/deviceRTLs/target_interface.h

+//===------------- target_interface.h - Target interfaces --------- C++ -*-===//
+//
+// 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 interfaces that must be implemented by each target.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _OMPTARGET_TARGET_INTERFACE_H_
+#define _OMPTARGET_TARGET_INTERFACE_H_
+
+#include "target_impl.h"
+
+// Calls to the NVPTX layer (assuming 1D layout)
+EXTERN int GetThreadIdInBlock();
+EXTERN int GetBlockIdInKernel();
+EXTERN int GetNumberOfBlocksInKernel();
+EXTERN int GetNumberOfThreadsInBlock();
+EXTERN unsigned GetWarpId();
+EXTERN unsigned GetLaneId();
+
+// Atomics
+extern DEVICE uint32_t __kmpc_atomic_add(uint32_t *, uint32_t);
+extern DEVICE uint32_t __kmpc_atomic_inc(uint32_t *, uint32_t);
+extern DEVICE uint32_t __kmpc_atomic_max(uint32_t *, uint32_t);
+extern DEVICE uint32_t __kmpc_atomic_exchange(uint32_t *, uint32_t);
+extern DEVICE uint32_t __kmpc_atomic_cas(uint32_t *, uint32_t, uint32_t);
+static_assert(sizeof(unsigned long long) == sizeof(uint64_t), "");
+extern DEVICE unsigned long long __kmpc_atomic_exchange(unsigned long long *,
+                                                        unsigned long long);
+extern DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *,
+                                                   unsigned long long);
+
+// Locks
+EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock);
+EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock);
+EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock);
+
+EXTERN void __kmpc_impl_threadfence();
+EXTERN void __kmpc_impl_threadfence_block();
+EXTERN void __kmpc_impl_threadfence_system();
+
+EXTERN double __kmpc_impl_get_wtick();
+EXTERN double __kmpc_impl_get_wtime();
+
+EXTERN void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi);
+EXTERN uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi);
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt();
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt();
+EXTERN uint32_t __kmpc_impl_smid();
+
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask();
+
+EXTERN int32_t __kmpc_impl_shfl_sync(__kmpc_impl_lanemask_t Mask, int32_t Var,
+                                     int32_t SrcLane);
+EXTERN int32_t __kmpc_impl_shfl_down_sync(__kmpc_impl_lanemask_t Mask,
+                                          int32_t Var, uint32_t Delta,
+                                          int32_t Width);
+
+EXTERN void __kmpc_impl_syncthreads();
+EXTERN void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask);
+
+// Kernel initialization
+EXTERN void __kmpc_impl_target_init();
+
+// Memory
+EXTERN void *__kmpc_impl_malloc(size_t);
+EXTERN void __kmpc_impl_free(void *);
+
+// Barrier until num_threads arrive.
+EXTERN void __kmpc_impl_named_sync(uint32_t num_threads);
+
+#endif // _OMPTARGET_TARGET_INTERFACE_H_