diff --git a/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt b/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt --- a/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt +++ b/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt @@ -16,7 +16,7 @@ # as of rocm-3.7, hsa is installed with cmake packages and kmt is found via hsa find_package(hsa-runtime64 QUIET 1.2.0 HINTS ${CMAKE_INSTALL_PREFIX} PATHS /opt/rocm) if (NOT ${hsa-runtime64_FOUND}) - libomptarget_say("Not building HSA plugin: hsa-runtime64 not found") + libomptarget_say("Not building AMDGPU plugin: hsa-runtime64 not found") return() endif() @@ -26,9 +26,15 @@ endif() if(NOT CMAKE_SYSTEM_PROCESSOR MATCHES "(x86_64)|(ppc64le)|(aarch64)$" AND CMAKE_SYSTEM_NAME MATCHES "Linux") - libomptarget_say("Not building amdgpu plugin: only support amdgpu in Linux x86_64, ppc64le, or aarch64 hosts.") + libomptarget_say("Not building AMDGPU plugin: only support AMDGPU in Linux x86_64, ppc64le, or aarch64 hosts.") return() endif() + +if (NOT LLVM_MAIN_INCLUDE_DIR) + libomptarget_say("Not building AMDGPU plugin: Missing definition for LLVM_MAIN_INCLUDE_DIR") + return() +endif() + libomptarget_say("Building amdgpu offloading plugin") ################################################################################ @@ -44,6 +50,7 @@ include_directories( ${CMAKE_CURRENT_SOURCE_DIR}/impl + ${LLVM_MAIN_INCLUDE_DIR} ) add_library(omptarget.rtl.amdgpu SHARED @@ -54,7 +61,7 @@ impl/system.cpp impl/utils.cpp impl/msgpack.cpp - src/rtl.cpp + src/rtl.cpp ) # Install plugin under the lib destination folder. diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp --- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp +++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp @@ -38,66 +38,7 @@ #include "Debug.h" #include "omptargetplugin.h" -// Get static gpu grid values from clang target-specific constants managed -// in the header file llvm/Frontend/OpenMP/OMPGridValues.h -// Copied verbatim to meet the requirement that libomptarget builds without -// a copy of llvm checked out nearby -namespace llvm { -namespace omp { -enum GVIDX { - /// The maximum number of workers in a kernel. - /// (THREAD_ABSOLUTE_LIMIT) - (GV_Warp_Size), might be issue for blockDim.z - GV_Threads, - /// The size reserved for data in a shared memory slot. - GV_Slot_Size, - /// The default value of maximum number of threads in a worker warp. - GV_Warp_Size, - /// Alternate warp size for some AMDGCN architectures. Same as GV_Warp_Size - /// for NVPTX. - GV_Warp_Size_32, - /// The number of bits required to represent the max number of threads in warp - GV_Warp_Size_Log2, - /// GV_Warp_Size * GV_Slot_Size, - GV_Warp_Slot_Size, - /// the maximum number of teams. - GV_Max_Teams, - /// Global Memory Alignment - GV_Mem_Align, - /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2)) - GV_Warp_Size_Log2_Mask, - // An alternative to the heavy data sharing infrastructure that uses global - // memory is one that uses device __shared__ memory. The amount of such space - // (in bytes) reserved by the OpenMP runtime is noted here. - GV_SimpleBufferSize, - // The absolute maximum team size for a working group - GV_Max_WG_Size, - // The default maximum team size for a working group - GV_Default_WG_Size, - // This is GV_Max_WG_Size / GV_WarpSize. 32 for NVPTX and 16 for AMDGCN. - GV_Max_Warp_Number, - /// The slot size that should be reserved for a working warp. - /// (~0u >> (GV_Warp_Size - GV_Warp_Size_Log2)) - GV_Warp_Size_Log2_MaskL -}; - -static constexpr unsigned AMDGPUGpuGridValues[] = { - 448, // GV_Threads - 256, // GV_Slot_Size - 64, // GV_Warp_Size - 32, // GV_Warp_Size_32 - 6, // GV_Warp_Size_Log2 - 64 * 256, // GV_Warp_Slot_Size - 128, // GV_Max_Teams - 256, // GV_Mem_Align - 63, // GV_Warp_Size_Log2_Mask - 896, // GV_SimpleBufferSize - 1024, // GV_Max_WG_Size, - 256, // GV_Defaut_WG_Size - 1024 / 64, // GV_Max_WG_Size / GV_WarpSize - 63 // GV_Warp_Size_Log2_MaskL -}; -} // namespace omp -} // namespace llvm +#include "llvm/Frontend/OpenMP/OMPGridValues.h" #ifndef TARGET_NAME #define TARGET_NAME AMDHSA