[Openmp-commits] [openmp] 7d2ecef - [openmp][libomptarget] Include header from LLVM source tree
via Openmp-commits
openmp-commits at lists.llvm.org
Thu Oct 15 07:46:51 PDT 2020
Author: JonChesterfield
Date: 2020-10-15T15:46:19+01:00
New Revision: 7d2ecef5ed11698ae106bfbf295c44d761c7f946
URL: https://github.com/llvm/llvm-project/commit/7d2ecef5ed11698ae106bfbf295c44d761c7f946
DIFF: https://github.com/llvm/llvm-project/commit/7d2ecef5ed11698ae106bfbf295c44d761c7f946.diff
LOG: [openmp][libomptarget] Include header from LLVM source tree
[openmp][libomptarget] Include header from LLVM source tree
The change is to the amdgpu plugin so is unlikely to break anything.
The point of contention is whether libomptarget can depend on LLVM.
A community discussion was cautiously not opposed yesterday.
This introduces a compile time dependency on the LLVM source tree, in this case
expressed as skipping the building of the plugin if LLVM_MAIN_INCLUDE_DIR is not
set. One the source files will #include llvm/Frontend/OpenMP/OMPGridValues.h,
instead of copy&pasting the numbers across.
For users that download the monorepo, the llvm tree is already on disk. This will
inconvenience users who download only the openmp source as a tar, as they would
now also have to download (at least a file or two) from the llvm source, if they want
to build the parts of the openmp project that (post this patch) depend on llvm.
There was interest expressed in going further - using llvm tools as part of
building libomp, or linking against llvm libraries. That seems less clear cut
an improvement and worthy of further discussion. This patch seeks only to change
policy to support openmp depending on the llvm source tree. Including in the
other direction, or using libraries / tools etc, are purposefully out of scope.
Reviewers are a best guess at interested parties, please feel free to add others
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D87841
Added:
Modified:
openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt b/openmp/libomptarget/plugins/amdgpu/CMakeLists.txt
index 7483e4e5c0ea..3882b777f5b1 100644
--- 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 @@ if(NOT LIBOMPTARGET_DEP_LIBELF_FOUND)
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 @@ endif()
include_directories(
${CMAKE_CURRENT_SOURCE_DIR}/impl
+ ${LLVM_MAIN_INCLUDE_DIR}
)
add_library(omptarget.rtl.amdgpu SHARED
@@ -54,7 +61,7 @@ add_library(omptarget.rtl.amdgpu SHARED
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
index 9ba27560d140..e0509a5f2b32 100644
--- 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
More information about the Openmp-commits
mailing list