[openmp][libomptarget] Include header from LLVM source tree

[JonChesterfield]

Summary

[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

[JonChesterfield]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

3–5

Somewhat unfortunately for this change, this header was added with the old license. Correcting here. Introduced by @saiislam and written by @gregrodgers from my team at AMD.

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

41

Lines on the left are byte for byte identical to the middle of this header, as that's where they were copied from. At some point in the future they're likely to accidentally diverge and I don't want to find that as the root cause while debugging.

[lebedev.ri]

As openmp and llvm are now developed in the same monorepo, the only uses this
dependency can break are those that download the monorepo, then delete the llvm
directory before calling cmake. That is not a compelling use case.

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

[JonChesterfield]

In D87841#2279743, @lebedev.ri wrote:

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

Thanks! I was not aware of that, will correct the description.

[saiislam]

AFAIK, all files which are required/should be available in libomptarget reside in llvm/Frontend/OpenMP/* , so may be an option can be provided for only-tarball users to get files only from this directory.
This kind of approach will also enforce people exporting stuff from llvm to libomptarget, to put everything in one place.

[saiislam]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

33

You might want to update these comments also. I have a NFC patch under review for exactly this D86232. Feel free to bring it in here.

[jdoerfert]

In D87841#2279743, @lebedev.ri wrote:

As openmp and llvm are now developed in the same monorepo, the only uses this
dependency can break are those that download the monorepo, then delete the llvm
directory before calling cmake. That is not a compelling use case.

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

Right. You can download the source for the sub-projects.
Though, others, e.g., clang, are already not compileable in isolation.

I see three options here:

1. Continue to copy code from llvm-core to openmp and back. This is "good" for people that compile libomp in isolation and "bad" for developers. Arguably, it is also "bad" for people that use libomp with clang because there is an increased risk we have an unnoticed mismatch between the two. This is not theoretical but already happened.
2. Copy the code from llvm-core to openmp during packaging time. This is a (small?) burden on the package maintainers but has all the advantages for now. However, it will get more complicated the more llvm-core features we want to use, e.g., tablegen.
3. Require users to download llvm-core and point to it. This is a small burden on the users but has all the advantages now and in the future.

This patch, on its own, picks 3). I'm happy with that.

[JonChesterfield]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

33

Thanks for the cross reference. I want to keep this review focused on accepting the new dependency edge.

[grokos]

In D87841#2280009, @jdoerfert wrote:

In D87841#2279743, @lebedev.ri wrote:

As openmp and llvm are now developed in the same monorepo, the only uses this
dependency can break are those that download the monorepo, then delete the llvm
directory before calling cmake. That is not a compelling use case.

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

Right. You can download the source for the sub-projects.
Though, others, e.g., clang, are already not compileable in isolation.

I see three options here:

1. Continue to copy code from llvm-core to openmp and back. This is "good" for people that compile libomp in isolation and "bad" for developers. Arguably, it is also "bad" for people that use libomp with clang because there is an increased risk we have an unnoticed mismatch between the two. This is not theoretical but already happened.
2. Copy the code from llvm-core to openmp during packaging time. This is a (small?) burden on the package maintainers but has all the advantages for now. However, it will get more complicated the more llvm-core features we want to use, e.g., tablegen.
3. Require users to download llvm-core and point to it. This is a small burden on the users but has all the advantages now and in the future.

This patch, on its own, picks 3). I'm happy with that.

4. At CMake time, check whether llvm-core is there; if yes use OMPGridValues.h from llvm-core, if no then download the header from some suitable location automatically. This eliminates the need for (2) and makes both proponents of (1) and proponents of (3) happy.

[lebedev.ri]

In D87841#2281494, @grokos wrote:

In D87841#2280009, @jdoerfert wrote:

In D87841#2279743, @lebedev.ri wrote:

As openmp and llvm are now developed in the same monorepo, the only uses this
dependency can break are those that download the monorepo, then delete the llvm
directory before calling cmake. That is not a compelling use case.

That is factually not accurate.
There are standalone tarballs, including one for omp:
https://releases.llvm.org/download.html#10.0.1 OpenMP Source code (.sig)

Right. You can download the source for the sub-projects.
Though, others, e.g., clang, are already not compileable in isolation.

I see three options here:

1. Continue to copy code from llvm-core to openmp and back. This is "good" for people that compile libomp in isolation and "bad" for developers. Arguably, it is also "bad" for people that use libomp with clang because there is an increased risk we have an unnoticed mismatch between the two. This is not theoretical but already happened.
2. Copy the code from llvm-core to openmp during packaging time. This is a (small?) burden on the package maintainers but has all the advantages for now. However, it will get more complicated the more llvm-core features we want to use, e.g., tablegen.
3. Require users to download llvm-core and point to it. This is a small burden on the users but has all the advantages now and in the future.

This patch, on its own, picks 3). I'm happy with that.

4. At CMake time, check whether llvm-core is there; if yes use OMPGridValues.h from llvm-core, if no then download the header from some suitable location automatically. This eliminates the need for (2) and makes both proponents of (1) and proponents of (3) happy.

*shrudder* Not that please.

[jdoerfert]

In D87841#2281494, @grokos wrote:

4. At CMake time, check whether llvm-core is there; if yes use OMPGridValues.h from llvm-core, if no then download the header from some suitable location automatically. This eliminates the need for (2) and makes both proponents of (1) and proponents of (3) happy.

Most of our bigger machine have no internet. Download after the fact might also confuse people. I think telling the user to download and point to instead, as we do with other dependences, is in the grant scheme of things less surprising.

[tianshilei1992]

In D87841#2281974, @jdoerfert wrote:

In D87841#2281494, @grokos wrote:

4. At CMake time, check whether llvm-core is there; if yes use OMPGridValues.h from llvm-core, if no then download the header from some suitable location automatically. This eliminates the need for (2) and makes both proponents of (1) and proponents of (3) happy.

Most of our bigger machine have no internet. Download after the fact might also confuse people. I think telling the user to download and point to instead, as we do with other dependences, is in the grant scheme of things less surprising.

I think it might be better to provide both options: if users don't specify anything, using online download by default. If users specify the path to the headers, it then does not download.

[JonChesterfield]

I strongly dislike build scripts pulling files off the internet. It's bad for machines that aren't connected to the web and upsets people in security

[fghanim]

I have no idea if this has been considered before or not, or if there is some technical/organizational difficulties against this option, but I think an option would be to create somewhere in the LLVM umbrella project for common Headers/Def.s between different LLVM subprojects (directory called "common" maybe?), which should help maintain independence between the projects without duplication, package maintainers will know to always include that in any package without having to worry about looking for shared files on a file by file bases, and users should always download that with whichever llvm package they are going to use.

FWIW.

[ABataev]

In D87841#2282280, @fghanim wrote:

I have no idea if this has been considered before or not, or if there is some technical/organizational difficulties against this option, but I think an option would be to create somewhere in the LLVM umbrella project for common Headers/Def.s between different LLVM subprojects (directory called "common" maybe?), which should help maintain independence between the projects without duplication, package maintainers will know to always include that in any package without having to worry about looking for shared files on a file by file bases, and users should always download that with whichever llvm package they are going to use.

FWIW.

+1 for this solution.

[JonChesterfield]

An additional repo which llvm, openmp, clang etc implicitly depend on would also work. It's roughly equivalent to treating llvm itself as that root dependency.

I suppose I view this set of repositories (now merged as the monorepo) as llvm + stuff, so it seems strange to require components in stuff to work wholly independently of llvm.

edit: A strong advantage to including from within llvm, instead of setting up a new top level directory to put shared things in, is that we can fix this maintenance hazard in openmp without waiting to reach consensus from the wider community on what to name, and what to put in, said new directory.

[JonChesterfield]

Ping. Can we have the include from llvm now, and postpone choosing a new top level directory until we have a longer list of things we'd like in it?

[JonChesterfield]

• rebase

[jdoerfert]

Can we move the cmake stuff in the main openmptarget logic so we can use llvm in all of openmp target.

[JonChesterfield]

In D87841#2330762, @jdoerfert wrote:

Can we move the cmake stuff in the main openmptarget logic so we can use llvm in all of openmp target.

Sounds great. Proposed as D89426, with the intent of landing that one first.

[jdoerfert]

LGTM, one comment.

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

6

Commit this separately as NFC

[saiislam]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

6

I will take care of it in my related NFC patch D86232.

[JonChesterfield]

-rebase

[JonChesterfield]

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

6

Done

llvm/include/llvm/Frontend/OpenMP/OMPGridValues.h

 //====--- OMPGridValues.h - Language-specific address spaces --*- C++ -*-====//
 //
-//                     The LLVM Compiler Infrastructure
-//
-// This file is distributed under the University of Illinois Open Source
-// License. See LICENSE.TXT for details.
+// 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
 //
 //===----------------------------------------------------------------------===//
 ///
 /// \file
 /// \brief Provides definitions for Target specific Grid Values
 ///
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_OPENMP_GRIDVALUES_H

 ///        multi-target binary. Both amdgcn and nvptx values are stored in
 ///        this file. In the future, should there be differences between GPUs
 ///        of the same architecture, then simply make a different array and
 ///        use the new array name.
 ///
 /// Example usage in clang:
 ///   const unsigned slot_size = ctx.GetTargetInfo().getGridValue(GV_Warp_Size);
 ///
 /// Example usage in libomptarget/deviceRTLs:
 ///   #include "OMPGridValues.h"
 ///   #ifdef __AMDGPU__
 ///     #define GRIDVAL AMDGPUGpuGridValues
 ///   #else
 ///     #define GRIDVAL NVPTXGpuGridValues
 ///   #endif
 ///   ... Then use this reference for GV_Warp_Size in the deviceRTL source.
 ///   GRIDVAL[GV_Warp_Size]

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

 #
 ##===----------------------------------------------------------------------===##
 
 ################################################################################
 
 # 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()
 
 if(NOT LIBOMPTARGET_DEP_LIBELF_FOUND)
   libomptarget_say("Not building AMDGPU plugin: LIBELF not found")
   return()
 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")
 
 ################################################################################
 # Define the suffix for the runtime messaging dumps.
 add_definitions(-DTARGET_NAME=AMDGPU)
 if(CMAKE_SYSTEM_PROCESSOR MATCHES "(ppc64le)|(aarch64)$")
    add_definitions(-DLITTLEENDIAN_CPU=1)
 endif()
 
 if(CMAKE_BUILD_TYPE MATCHES Debug)
   add_definitions(-DDEBUG)
 endif()
 
 include_directories(
   ${CMAKE_CURRENT_SOURCE_DIR}/impl
+  ${LLVM_MAIN_INCLUDE_DIR}
 )
 
 add_library(omptarget.rtl.amdgpu SHARED
       impl/atmi.cpp
       impl/atmi_interop_hsa.cpp
       impl/data.cpp
       impl/machine.cpp
       impl/system.cpp
       impl/utils.cpp
       impl/msgpack.cpp
-      src/rtl.cpp 
+      src/rtl.cpp
       )
 
 # Install plugin under the lib destination folder.
 # When we build for debug, OPENMP_LIBDIR_SUFFIX get set to -debug
 install(TARGETS omptarget.rtl.amdgpu LIBRARY DESTINATION "lib${OPENMP_LIBDIR_SUFFIX}")
 
 set_property(TARGET omptarget.rtl.amdgpu PROPERTY INSTALL_RPATH "$ORIGIN")
 target_link_libraries(

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <dlfcn.h>
 #include <elf.h>
 #include <ffi.h>
 #include <fstream>
 #include <iostream>
 #include <libelf.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
 #include <thread>
 #include <unordered_map>
 #include <vector>
 
 // Header from ATMI interface
 #include "atmi_interop_hsa.h"
 #include "atmi_runtime.h"
 
 #include "internal.h"
 
 #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
 #endif
 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 
 int print_kernel_trace;