From dd8a7fcdd73dd63529b81bf9f72c7529dfe99ec3 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Fri, 13 Dec 2019 16:24:10 -0500
Subject: [PATCH] Revert "[libomptarget] Move resource id functions into target
 specific code, implement for amdgcn"

This reverts commit dbb3fec8adfc4ac3fbf31f51f294427dbabbebb2 since it
breaks the NVPTX tests.
---
 openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h | 8 --------
 openmp/libomptarget/deviceRTLs/common/src/support.cu    | 8 ++++++++
 openmp/libomptarget/deviceRTLs/common/support.h         | 6 +++++-
 openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h  | 6 ------
 4 files changed, 13 insertions(+), 15 deletions(-)

diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
index 5082d46..62cbfb0 100644
--- a/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.h
@@ -136,14 +136,6 @@ 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 ea1fc38..2f992f2 100644
--- a/openmp/libomptarget/deviceRTLs/common/src/support.cu
+++ b/openmp/libomptarget/deviceRTLs/common/src/support.cu
@@ -98,6 +98,14 @@ 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 400d264..8cffd91 100644
--- a/openmp/libomptarget/deviceRTLs/common/support.h
+++ b/openmp/libomptarget/deviceRTLs/common/support.h
@@ -1,4 +1,4 @@
-//===--------- support.h - OpenMP GPU support functions ---------- CUDA -*-===//
+//===--------- support.h - NVPTX OpenMP 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,6 +51,10 @@ 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 161cd6c..fe36a46 100644
--- a/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.h
@@ -167,10 +167,4 @@ 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
-- 
2.7.4