diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
--- a/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/CMakeLists.txt
@@ -54,7 +54,8 @@
   ${CMAKE_CURRENT_SOURCE_DIR}
   DIRECTORY)
 
-set(cuda_sources
+set(rtl_sources
+  ${CMAKE_CURRENT_SOURCE_DIR}/src/resource_id.hip
   ${devicertl_base_directory}/common/src/cancel.cu
   ${devicertl_base_directory}/common/src/critical.cu)
 
@@ -63,6 +64,7 @@
   ${CMAKE_CURRENT_SOURCE_DIR}/src/target_impl.h
   ${devicertl_base_directory}/common/debug.h
   ${devicertl_base_directory}/common/device_environment.h
+  ${devicertl_base_directory}/common/resource_id.h
   ${devicertl_base_directory}/common/state-queue.h
   ${devicertl_base_directory}/common/state-queuei.h
   ${devicertl_base_directory}/common/support.h)
@@ -80,7 +82,7 @@
   set(mcpus ${LIBOMPTARGET_AMDGCN_GFXLIST})
 endif()
 
-macro(add_cuda_bc_library)
+macro(add_rtl_bc_library)
   set(cu_cmd ${AOMP_BINDIR}/clang++
     -std=c++11
     -fcuda-rdc
@@ -120,7 +122,7 @@
 
 foreach(mcpu ${mcpus})
   set(bc_files)
-  add_cuda_bc_library(${cuda_sources})
+  add_rtl_bc_library(${rtl_sources})
 
   set(bc_libname lib${libname}-${mcpu}.bc)
   add_custom_command(
diff --git a/openmp/libomptarget/deviceRTLs/amdgcn/src/resource_id.hip b/openmp/libomptarget/deviceRTLs/amdgcn/src/resource_id.hip
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/amdgcn/src/resource_id.hip
@@ -0,0 +1,25 @@
+//===--- resource_id.hip - AMDGCN OpenMP resource id functions --- HIP -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Accessors for low level ids of resources
+//
+//===----------------------------------------------------------------------===//
+
+#include "common/resource_id.h"
+
+// 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);
+
+DEVICE int GetThreadIdInBlock() { return __builtin_amdgcn_workitem_id_x(); }
+
+DEVICE int GetBlockIdInKernel() { return __builtin_amdgcn_workgroup_id_x(); }
+
+DEVICE int GetNumberOfBlocksInKernel() { return __ockl_get_num_groups(0); }
+
+DEVICE int GetNumberOfThreadsInBlock() { return __ockl_get_local_size(0); }
diff --git a/openmp/libomptarget/deviceRTLs/common/resource_id.h b/openmp/libomptarget/deviceRTLs/common/resource_id.h
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/common/resource_id.h
@@ -0,0 +1,27 @@
+//===--- resource_id.h - OpenMP GPU resource id 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Accessors for low level ids of resources
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OMPTARGET_RESOURCE_ID_H
+#define OMPTARGET_RESOURCE_ID_H
+
+#include "target_impl.h"
+
+DEVICE int GetThreadIdInBlock();
+DEVICE int GetBlockIdInKernel();
+DEVICE int GetNumberOfBlocksInKernel();
+DEVICE int GetNumberOfThreadsInBlock();
+DEVICE unsigned GetWarpId();
+DEVICE unsigned GetLaneId();
+
+#endif
+
+
diff --git a/openmp/libomptarget/deviceRTLs/common/src/support.cu b/openmp/libomptarget/deviceRTLs/common/src/support.cu
--- a/openmp/libomptarget/deviceRTLs/common/src/support.cu
+++ b/openmp/libomptarget/deviceRTLs/common/src/support.cu
@@ -98,14 +98,6 @@
 //
 ////////////////////////////////////////////////////////////////////////////////
 
-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
--- 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.
@@ -15,6 +15,7 @@
 
 #include "interface.h"
 #include "target_impl.h"
+#include "resource_id.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
@@ -51,10 +52,6 @@
 ////////////////////////////////////////////////////////////////////////////////
 
 // 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/CMakeLists.txt b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
--- a/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
+++ b/openmp/libomptarget/deviceRTLs/nvptx/CMakeLists.txt
@@ -59,6 +59,7 @@
       ${devicertl_common_directory}/src/omptarget.cu
       ${devicertl_common_directory}/src/parallel.cu
       src/reduction.cu
+      src/resource_id.cu
       ${devicertl_common_directory}/src/support.cu
       ${devicertl_common_directory}/src/sync.cu
       ${devicertl_common_directory}/src/task.cu
diff --git a/openmp/libomptarget/deviceRTLs/nvptx/src/resource_id.cu b/openmp/libomptarget/deviceRTLs/nvptx/src/resource_id.cu
new file mode 100644
--- /dev/null
+++ b/openmp/libomptarget/deviceRTLs/nvptx/src/resource_id.cu
@@ -0,0 +1,23 @@
+//===--- resource_id.cu - NVPTX OpenMP resource id 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Accessors for low level ids of resources
+//
+//===----------------------------------------------------------------------===//
+
+#include "common/resource_id.h"
+
+// 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; }