[libomptarget] Fix devicertl build

[JonChesterfield]

Summary

[libomptarget] Fix devicertl build

The target specific functions in target_interface are extern C, but the
implementations for nvptx were mostly C++ mangling. That worked out as
a quirk of DEVICE macro expanding to nothing, except for shuffle.h which
only forward declared the functions with C++ linkage.

Also implements GetWarpSize, as used by shuffle, and includes target_interface
in nvptx target_impl.cu to help catch future divergence between interface and
implementation.

[JonChesterfield]

• drop accidental amdgcn part

[JonChesterfield]

openmp/libomptarget/deviceRTLs/target_interface.h

28–32

These are 'extern' by default, but not 'EXTERN' which is a macro that expands to extern "C", which in this context seems confusing.

[jdoerfert]

Assuming this unblocks the build, LG. We should clean up the extern mess though.

[JonChesterfield]

Yep. Also, DEVICE now always expands to whitespace, so that can go.

[tianshilei1992]

In D98651#2627076, @JonChesterfield wrote:

Yep. Also, DEVICE now always expands to whitespace, so that can go.

Code for AMD platform is the only reason we still have DEVICE in common parts. ;-)

[JonChesterfield]

In D98651#2627106, @tianshilei1992 wrote:

In D98651#2627076, @JonChesterfield wrote:

Yep. Also, DEVICE now always expands to whitespace, so that can go.

Code for AMD platform is the only reason we still have DEVICE in common parts. ;-)

Your info is out of date, HIP build is no more. Patched at D98655

openmp/libomptarget/deviceRTLs/CMakeLists.txt

 ##===----------------------------------------------------------------------===##
 #
 # 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
 #
 # ##===----------------------------------------------------------------------===##
 #
 # Build a device RTL for each available machine.
 #
 ##===----------------------------------------------------------------------===##
 
+add_subdirectory(amdgcn)
 add_subdirectory(nvptx)

openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt

foreach(mcpu ${mcpus})
   add_custom_command(
     OUTPUT ${bc_libname}
     COMMAND ${AOMP_BINDIR}/llvm-link ${bc_files} | ${AOMP_BINDIR}/opt --always-inline -o ${OUTPUTDIR}/${bc_libname}
     DEPENDS ${bc_files})
 
   add_custom_target(lib${libname}-${mcpu} ALL DEPENDS ${bc_libname})
 
   install(FILES ${OUTPUTDIR}/${bc_libname}
-     DESTINATION "${OPENMP_INSTALL_LIBDIR}/libdevice"
+     DESTINATION "${OPENMP_INSTALL_LIBDIR}"
   )
 endforeach()

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

 
 EXTERN int GetNumberOfThreadsInBlock() {
   return get_workgroup_dim(__builtin_amdgcn_workgroup_id_x(),
                            __builtin_amdgcn_grid_size_x(),
                            __builtin_amdgcn_workgroup_size_x());
 }
 
 EXTERN unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+EXTERN unsigned GetWarpSize() { return WARPSIZE; }
 EXTERN unsigned GetLaneId() {
   return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
 }
 
 EXTERN uint32_t __kmpc_amdgcn_gpu_num_threads() {
   return GetNumberOfThreadsInBlock();
 }
 

openmp/libomptarget/deviceRTLs/common/include/target/shuffle.h

 int64_t __kmpc_shuffle_int64(int64_t val, int16_t delta, int16_t size);
 }
 
 ///}
 
 /// Forward declarations
 ///
 ///{
+extern "C" {
 unsigned GetLaneId();
 unsigned GetWarpSize();
 void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi);
 uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi);
+}
 ///}
 
 /// Fallback implementations of the shuffle sync idiom.
 /// Unavailable at present (would error at link time if used).
 ///
 ///{
 
 int32_t __kmpc_impl_shfl_sync(uint64_t Mask, int32_t Var, int32_t SrcLane);
 
 int32_t __kmpc_impl_shfl_down_sync(uint64_t Mask, int32_t Var, uint32_t Delta,
                                    int32_t Width);
 
 ///}
 
 /// AMDGCN implementations of the shuffle sync idiom.
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
 inline int32_t __kmpc_impl_shfl_sync(uint64_t Mask, int32_t Var,
                                      int32_t SrcLane) {
   int Width = GetWarpSize();
   int Self = GetLaneId();
   int Index = SrcLane + (Self & ~(Width - 1));
   return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
 }
 
 inline int32_t __kmpc_impl_shfl_down_sync(uint64_t Mask, 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);
 }
 
 #pragma omp end declare variant
 ///}
 
 /// NVPTX implementations of the shuffle and shuffle sync idiom.
 ///
 ///{
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 inline int32_t __kmpc_impl_shfl_sync(uint64_t Mask, int32_t Var,
                                      int32_t SrcLane) {
   return __nvvm_shfl_idx_i32(Var, SrcLane, 0x1f);
 }
 
 inline int32_t __kmpc_impl_shfl_down_sync(uint64_t Mask, int32_t Var,
                                           uint32_t Delta, int32_t Width) {
   int32_t T = ((GetWarpSize() - Width) << 8) | 0x1f;
   return __nvvm_shfl_down_i32(Var, Delta, T);
 }
 
 #pragma omp end declare variant
 ///}
 
 #pragma omp end declare target
 
 #endif

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

 //===---------- target_impl.cu - NVPTX 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
 //
 //===----------------------------------------------------------------------===//
 #pragma omp declare target
 
 #include "common/debug.h"
 #include "target_impl.h"
+#include "target_interface.h"
 
-DEVICE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
+EXTERN void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "l"(val));
 }
 
-DEVICE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
+EXTERN uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   uint64_t val;
   asm volatile("mov.b64 %0, {%1,%2};" : "=l"(val) : "r"(lo), "r"(hi));
   return val;
 }
 
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_lanemask_lt() {
   __kmpc_impl_lanemask_t res;
   asm("mov.u32 %0, %%lanemask_lt;" : "=r"(res));
   return res;
 }
 
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_lanemask_gt() {
   __kmpc_impl_lanemask_t res;
   asm("mov.u32 %0, %%lanemask_gt;" : "=r"(res));
   return res;
 }
 
-DEVICE uint32_t __kmpc_impl_smid() {
+EXTERN uint32_t __kmpc_impl_smid() {
   uint32_t id;
   asm("mov.u32 %0, %%smid;" : "=r"(id));
   return id;
 }
 
-DEVICE double __kmpc_impl_get_wtick() {
+EXTERN double __kmpc_impl_get_wtick() {
   // Timer precision is 1ns
   return ((double)1E-9);
 }
 
-DEVICE double __kmpc_impl_get_wtime() {
+EXTERN double __kmpc_impl_get_wtime() {
   unsigned long long nsecs;
   asm("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
   return (double)nsecs * __kmpc_impl_get_wtick();
 }
 
-DEVICE __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
+EXTERN __kmpc_impl_lanemask_t __kmpc_impl_activemask() {
   unsigned int Mask;
   asm volatile("activemask.b32 %0;" : "=r"(Mask));
   return Mask;
 }
 
-DEVICE void __kmpc_impl_syncthreads() { __syncthreads(); }
+EXTERN void __kmpc_impl_syncthreads() { __syncthreads(); }
 
-DEVICE void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
+EXTERN void __kmpc_impl_syncwarp(__kmpc_impl_lanemask_t Mask) {
   __nvvm_bar_warp_sync(Mask);
 }
 
 // NVPTX specific kernel initialization
-DEVICE void __kmpc_impl_target_init() { /* nvptx needs no extra setup */
+EXTERN void __kmpc_impl_target_init() { /* nvptx needs no extra setup */
 }
 
 // Barrier until num_threads arrive.
-DEVICE void __kmpc_impl_named_sync(uint32_t num_threads) {
+EXTERN void __kmpc_impl_named_sync(uint32_t num_threads) {
   // The named barrier for active parallel threads of a team in an L1 parallel
   // region to synchronize with each other.
   int barrier = 1;
   asm volatile("bar.sync %0, %1;"
                :
                : "r"(barrier), "r"(num_threads)
                : "memory");
 }
 
-DEVICE void __kmpc_impl_threadfence() { __nvvm_membar_gl(); }
-DEVICE void __kmpc_impl_threadfence_block() { __nvvm_membar_cta(); }
-DEVICE void __kmpc_impl_threadfence_system() { __nvvm_membar_sys(); }
+EXTERN void __kmpc_impl_threadfence() { __nvvm_membar_gl(); }
+EXTERN void __kmpc_impl_threadfence_block() { __nvvm_membar_cta(); }
+EXTERN void __kmpc_impl_threadfence_system() { __nvvm_membar_sys(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
-DEVICE int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
-DEVICE int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
-DEVICE int GetNumberOfBlocksInKernel() {
+EXTERN int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
+EXTERN int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
+EXTERN int GetNumberOfBlocksInKernel() {
   return __nvvm_read_ptx_sreg_nctaid_x();
 }
-DEVICE int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
-DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
-DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
+EXTERN int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
+EXTERN unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
+EXTERN unsigned GetWarpSize() { return WARPSIZE; }
+EXTERN unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 
 // Atomics
 DEVICE uint32_t __kmpc_atomic_add(uint32_t *Address, uint32_t Val) {
   return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 DEVICE uint32_t __kmpc_atomic_inc(uint32_t *Address, uint32_t Val) {
   return __nvvm_atom_inc_gen_ui(Address, Val);
 }
DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *Address,
                                             unsigned long long Val) {
   return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 
 #define __OMP_SPIN 1000
 #define UNSET 0u
 #define SET 1u
 
-DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock) {
+EXTERN void __kmpc_impl_init_lock(omp_lock_t *lock) {
   __kmpc_impl_unset_lock(lock);
 }
 
-DEVICE void __kmpc_impl_destroy_lock(omp_lock_t *lock) {
+EXTERN void __kmpc_impl_destroy_lock(omp_lock_t *lock) {
   __kmpc_impl_unset_lock(lock);
 }
 
-DEVICE void __kmpc_impl_set_lock(omp_lock_t *lock) {
+EXTERN void __kmpc_impl_set_lock(omp_lock_t *lock) {
   // TODO: not sure spinning is a good idea here..
   while (__kmpc_atomic_cas(lock, UNSET, SET) != UNSET) {
     int32_t start = __nvvm_read_ptx_sreg_clock();
     int32_t now;
     for (;;) {
       now = __nvvm_read_ptx_sreg_clock();
       int32_t cycles = now > start ? now - start : now + (0xffffffff - start);
       if (cycles >= __OMP_SPIN * GetBlockIdInKernel()) {
         break;
       }
     }
   } // wait for 0 to be the read value
 }
 
-DEVICE void __kmpc_impl_unset_lock(omp_lock_t *lock) {
+EXTERN void __kmpc_impl_unset_lock(omp_lock_t *lock) {
   (void)__kmpc_atomic_exchange(lock, UNSET);
 }
 
-DEVICE int __kmpc_impl_test_lock(omp_lock_t *lock) {
+EXTERN int __kmpc_impl_test_lock(omp_lock_t *lock) {
   return __kmpc_atomic_add(lock, 0u);
 }
 
-DEVICE void *__kmpc_impl_malloc(size_t x) { return malloc(x); }
-DEVICE void __kmpc_impl_free(void *x) { free(x); }
+EXTERN void *__kmpc_impl_malloc(size_t x) { return malloc(x); }
+EXTERN void __kmpc_impl_free(void *x) { free(x); }
 
 #pragma omp end declare target

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 GetWarpSize();
 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);
+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), "");
-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);
+DEVICE unsigned long long __kmpc_atomic_exchange(unsigned long long *,
+DEVICE unsigned long long __kmpc_atomic_add(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);