maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

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

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

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: jdoerfert

Subscribers: jvesely, mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71382
This commit is contained in:
Jon Chesterfield 2019-12-12 22:49:01 +00:00
parent 259a9b1039
commit dbb3fec8ad
4 changed files with 15 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
View File

@ -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

8
openmp/libomptarget/deviceRTLs/common/src/support.cu
View File

@ -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); }

6
openmp/libomptarget/deviceRTLs/common/support.h
View File

@ -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();

6
openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
View File

@ -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