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

[JonChesterfield]

Summary

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

[JonChesterfield]

Had some trouble picking a name for this file. It's a place to put the four accesses to magic cuda variables so that non-cuda targets can provide an equivalent, with its own header as that seemed better than implementing support.h across multiple sources.

[JonChesterfield]

• rebase

[ABataev]

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

1 ↗

(On Diff #233537)

We have target-specific files already, like target_impl.h, maybe just put all these stuff in this file?

[JonChesterfield]

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

1 ↗

(On Diff #233537)

target_impl is becoming a bit of a dumping ground, but sure - another four functions won't hurt

[JonChesterfield]

• Fold resource_id into target_impl as requested

[JonChesterfield]

• remove whitespace from comment

[jdoerfert]

LGTM.

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