diff options
4 files changed, 15 insertions, 13 deletions
diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h index 192e355c62f..823fc283dc8 100644 --- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h @@ -130,6 +130,14 @@ 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 diff --git a/openmp/libomptarget/deviceRTLs/common/src/support.cu b/openmp/libomptarget/deviceRTLs/common/src/support.cu index 2f992f2778e..ea1fc3841ad 100644 --- a/openmp/libomptarget/deviceRTLs/common/src/support.cu +++ b/openmp/libomptarget/deviceRTLs/common/src/support.cu @@ -98,14 +98,6 @@ DEVICE bool checkRuntimeInitialized(kmp_Ident *loc) { // //////////////////////////////////////////////////////////////////////////////// -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); } diff --git a/openmp/libomptarget/deviceRTLs/common/support.h b/openmp/libomptarget/deviceRTLs/common/support.h index 8cffd91c9f3..400d2649afd 100644 --- a/openmp/libomptarget/deviceRTLs/common/support.h +++ b/openmp/libomptarget/deviceRTLs/common/support.h @@ -1,4 +1,4 @@ -//===--------- 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. @@ -51,10 +51,6 @@ DEVICE bool checkRuntimeInitialized(kmp_Ident *loc); //////////////////////////////////////////////////////////////////////////////// // get low level ids of resources -DEVICE int GetThreadIdInBlock(); -DEVICE int GetBlockIdInKernel(); -DEVICE int GetNumberOfBlocksInKernel(); -DEVICE int GetNumberOfThreadsInBlock(); DEVICE unsigned GetWarpId(); DEVICE unsigned GetLaneId(); diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h index fe36a46c5cd..161cd6cac11 100644 --- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h +++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h @@ -167,4 +167,10 @@ 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 |
