[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

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

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/CMakeLists.txt

 if(LIBOMPTARGET_NVPTX_DEBUG)
   set(CUDA_DEBUG -DOMPTARGET_NVPTX_DEBUG=-1)
 endif()
 
 get_filename_component(devicertl_base_directory
   ${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)
 
 set(h_files
   ${CMAKE_CURRENT_SOURCE_DIR}/src/amdgcn_interface.h
   ${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)
 
 # for both in-tree and out-of-tree build
 if (NOT CMAKE_ARCHIVE_OUTPUT_DIRECTORY)
   set(OUTPUTDIR ${CMAKE_CURRENT_BINARY_DIR})
 else()
   set(OUTPUTDIR ${CMAKE_ARCHIVE_OUTPUT_DIRECTORY})
 endif()
 
 # create libraries
 set(mcpus gfx700 gfx701 gfx801 gfx803 gfx900)
 if (DEFINED LIBOMPTARGET_AMDGCN_GFXLIST)
   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
     -fvisibility=default
     --cuda-device-only
     -Wno-unused-value
     -x hip
     -O${optimization_level}
macro(add_rtl_bc_library)
 
   list(APPEND bc_files linkout.cuda.${mcpu}.bc)
 endmacro()
 
 set(libname "omptarget-amdgcn")
 
 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(
     OUTPUT ${bc_libname}
     COMMAND ${AOMP_BINDIR}/llvm-link ${bc_files} | ${AOMP_BINDIR}/opt --always-inline -o ${OUTPUTDIR}/${bc_libname}
     DEPENDS ${bc_files})
 
   add_custom_target(lib${libname}-${mcpu} ALL DEPENDS ${bc_libname})
 
   install(FILES ${OUTPUTDIR}/${bc_libname}
      DESTINATION "${OPENMP_INSTALL_LIBDIR}/libdevice"
   )
 endforeach()

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

+//===--- 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); }

openmp/libomptarget/deviceRTLs/common/resource_id.h

+//===--- 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
+
+

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.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef OMPTARGET_SUPPORT_H
 #define OMPTARGET_SUPPORT_H
 
 #include "interface.h"
 #include "target_impl.h"
+#include "resource_id.h"
 
 ////////////////////////////////////////////////////////////////////////////////
 // Execution Parameters
 ////////////////////////////////////////////////////////////////////////////////
 enum ExecutionMode {
   Generic = 0x00u,
   Spmd = 0x01u,
   ModeMask = 0x01u,

 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/CMakeLists.txt

set(cuda_src_files
       ${devicertl_common_directory}/src/cancel.cu
       ${devicertl_common_directory}/src/critical.cu
       src/data_sharing.cu
       src/libcall.cu
       ${devicertl_common_directory}/src/loop.cu
       ${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
   )
 
   set(omp_data_objects src/omp_data.cu)
 
   # Get the compute capability the user requested or use SM_35 by default.

openmp/libomptarget/deviceRTLs/nvptx/src/resource_id.cu

+//===--- 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; }