Revert "Revert "[libomptarget] Move resource id functions into target specific code, implement for amdgcn""

[JonChesterfield]

Summary

This reverts commit dd8a7fcdd73dd63529b81bf9f72c7529dfe99ec3.

Alexey reports undefined symbols for the new inline functions defined in target_impl.h
This does not reproduce for me for nvptx, or amdgcn, under release or debug builds.

I believe the patch is fine, based on:

• the semantics of an inline function in C++ (the cuda INLINE functions end up as linkonce_odr in IR), which are only legal to drop if they have no uses
• the code generated from a debug build of clang 9 does not show these undef symbols
• the tests pass
• the code is trivial

To progress from here I either need:

• A tie break - someone to play the role of CI in determining whether the patch works
• Alexey to provide sufficient information about his build for me to reproduce the failure
• Alexey to debug why the symbols are disappearing for him and report back

[JonChesterfield]

One would get missing symbols if objects compiled against the previous support.h were linked to ones compiled against the files here, e.g. if one attempted an incremental build of deviceRTL despite the dependency graph changing between revisions.

[jdoerfert]

In D71502#1784692, @JonChesterfield wrote:

One would get missing symbols if objects compiled against the previous support.h were linked to ones compiled against the files here, e.g. if one attempted an incremental build of deviceRTL despite the dependency graph changing between revisions.

Let's wait till @ABataev comes back with the system setup that caused this and the names of the symbols that were missing.

[JonChesterfield]

Ping. I'm blocked until this is resolved.

[ABataev]

In D71502#1785962, @JonChesterfield wrote:

Ping. I'm blocked until this is resolved.

Yes, it works, sorry for the delay.

[JonChesterfield]

In D71502#1785967, @ABataev wrote:

In D71502#1785962, @JonChesterfield wrote:

Ping. I'm blocked until this is resolved.

Yes, it works, sorry for the delay.

That's a considerable relief, thanks. Will land this.

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

INLINE void __kmpc_impl_named_sync(int barrier, uint32_t num_threads) {
   // So it is ok to issue a full workgroup barrier here.
   __builtin_amdgcn_s_barrier();
 }
 
 EXTERN void __kmpc_impl_threadfence(void);
 EXTERN void __kmpc_impl_threadfence_block(void);
 EXTERN void __kmpc_impl_threadfence_system(void);
 
+// Calls to the AMDGCN layer (assuming 1D layout)
+EXTERN uint64_t __ockl_get_local_size(uint32_t);
+EXTERN uint64_t __ockl_get_num_groups(uint32_t);
+INLINE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
+INLINE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
+INLINE int GetNumberOfBlocksInKernel() { return __ockl_get_num_groups(0); }
+INLINE int GetNumberOfThreadsInBlock() { return __ockl_get_local_size(0); }
+
 // DEVICE versions of part of libc
 extern "C" {
 DEVICE __attribute__((noreturn)) void
 __assertfail(const char *, const char *, unsigned, const char *, size_t);
 INLINE static void __assert_fail(const char *__message, const char *__file,
                                  unsigned int __line, const char *__function) {
   __assertfail(__message, __file, __line, __function, sizeof(char));
 }
 DEVICE int printf(const char *, ...);
 }
 
 #endif

openmp/libomptarget/deviceRTLs/common/src/support.cu

 ////////////////////////////////////////////////////////////////////////////////
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // Calls to the NVPTX layer  (assuming 1D layout)
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-DEVICE int GetThreadIdInBlock() { return threadIdx.x; }
-
-DEVICE int GetBlockIdInKernel() { return blockIdx.x; }
-
-DEVICE int GetNumberOfBlocksInKernel() { return gridDim.x; }
-
-DEVICE int GetNumberOfThreadsInBlock() { return blockDim.x; }
-
 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 
 DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 
 ////////////////////////////////////////////////////////////////////////////////
 //
 // Calls to the Generic Scheme Implementation Layer (assuming 1D layout)
 //

openmp/libomptarget/deviceRTLs/common/support.h

-//===--------- support.h - NVPTX OpenMP support functions -------- CUDA -*-===//
+//===--------- support.h - OpenMP GPU support functions ---------- 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Wrapper to some functions natively supported by the GPU.

 DEVICE bool checkRuntimeUninitialized(kmp_Ident *loc);
 DEVICE bool checkRuntimeInitialized(kmp_Ident *loc);
 
 ////////////////////////////////////////////////////////////////////////////////
 // get info from machine
 ////////////////////////////////////////////////////////////////////////////////
 
 // get low level ids of resources
-DEVICE int GetThreadIdInBlock();
-DEVICE int GetBlockIdInKernel();
-DEVICE int GetNumberOfBlocksInKernel();
-DEVICE int GetNumberOfThreadsInBlock();
 DEVICE unsigned GetWarpId();
 DEVICE unsigned GetLaneId();
 
 // get global ids to locate tread/team info (constant regardless of OMP)
 DEVICE int GetLogicalThreadIdInBlock(bool isSPMDExecutionMode);
 DEVICE int GetMasterThreadID();
 DEVICE int GetNumberOfWorkersInTeam();
 

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

asm volatile("bar.sync %0, %1;"
                : "r"(barrier), "r"(num_threads)
                : "memory");
 }
 
 INLINE void __kmpc_impl_threadfence(void) { __threadfence(); }
 INLINE void __kmpc_impl_threadfence_block(void) { __threadfence_block(); }
 INLINE void __kmpc_impl_threadfence_system(void) { __threadfence_system(); }
 
+// Calls to the NVPTX layer (assuming 1D layout)
+INLINE int GetThreadIdInBlock() { return threadIdx.x; }
+INLINE int GetBlockIdInKernel() { return blockIdx.x; }
+INLINE int GetNumberOfBlocksInKernel() { return gridDim.x; }
+INLINE int GetNumberOfThreadsInBlock() { return blockDim.x; }
+
 #endif