[OpenMP][libomptarget] New plugin infrastructure and new CUDA plugin

This patch adds a new infrastructure for OpenMP target plugins. It also implements the CUDA and GenericELF64bit plugins under this new infrastructure. We place the sources in a separate directory named plugins-nextgen, and we build the new plugins as different plugin libraries. The original plugins, which remain untouched, will be used by default. However, the user can change this behavior at run-time through the boolean envar LIBOMPTARGET_NEXTGEN_PLUGINS. If enabled, the libomptarget will try to load the NextGen version of each plugin, falling back to the original if they are not present or valid.

The idea of this new plugin infrastructure is to implement the common parts of target plugins in generic classes (defined in files inside plugins-next/common/PluginInterface folder), and then, each specific plugin defines its own specific classes inheriting from the common ones. In this way, most logic remains on the common interface while reducing the plugin-specific source code. It is also beneficial in the sense that now most code and behavior are the same across the different plugins. As an example, we define classes for a plugin, a device, a device image, a stream manager, etc. The plugin object (a single instance per plugin library) holds different device objects (i.e., one per available device), while these latter are the responsible for managing its own resources.

Most code on this patch is based on the changes made by @jdoerfert (Johannes Doerfert)

Reviewed By: jhuber6, jdoerfert

Differential Revision: https://reviews.llvm.org/D134396

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

@@ -56,24 +56,24 @@ namespace omp {
 
 struct GV {
   /// The size reserved for data in a shared memory slot.
-  const unsigned GV_Slot_Size;
+  unsigned GV_Slot_Size;
   /// The default value of maximum number of threads in a worker warp.
-  const unsigned GV_Warp_Size;
+  unsigned GV_Warp_Size;
 
   constexpr unsigned warpSlotSize() const {
     return GV_Warp_Size * GV_Slot_Size;
   }
 
   /// the maximum number of teams.
-  const unsigned GV_Max_Teams;
+  unsigned GV_Max_Teams;
   // 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.
-  const unsigned GV_SimpleBufferSize;
+  unsigned GV_SimpleBufferSize;
   // The absolute maximum team size for a working group
-  const unsigned GV_Max_WG_Size;
+  unsigned GV_Max_WG_Size;
   // The default maximum team size for a working group
-  const unsigned GV_Default_WG_Size;
+  unsigned GV_Default_WG_Size;
 
   constexpr unsigned maxWarpNumber() const {
     return GV_Max_WG_Size / GV_Warp_Size;

openmp/libomptarget/CMakeLists.txt

@@ -85,6 +85,7 @@ set(LIBOMPTARGET_LLVM_LIBRARY_DIR "${LLVM_LIBRARY_DIR}" CACHE STRING
 
 # Build offloading plugins and device RTLs if they are available.
 add_subdirectory(plugins)
+add_subdirectory(plugins-nextgen)
 add_subdirectory(DeviceRTL)
 add_subdirectory(tools)
 

openmp/libomptarget/include/Utilities.h

@@ -0,0 +1,200 @@
+//===------- Utilities.h - Target independent OpenMP target RTL -- C++ ----===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Routines and classes used to provide useful functionalities like string
+// parsing and environment variables.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef OPENMP_LIBOMPTARGET_INCLUDE_UTILITIES_H
+#define OPENMP_LIBOMPTARGET_INCLUDE_UTILITIES_H
+
+#include "llvm/ADT/STLFunctionalExtras.h"
+
+#include "Debug.h"
+
+#include <algorithm>
+#include <cassert>
+#include <cstdint>
+#include <cstdlib>
+#include <functional>
+#include <sstream>
+#include <string>
+
+namespace llvm {
+namespace omp {
+namespace target {
+
+/// Utility class for parsing strings to other types.
+struct StringParser {
+  /// Parse a string to another type.
+  template <typename Ty> static bool parse(const char *Value, Ty &Result);
+};
+
+/// Class for reading and checking environment variables. Currently working with
+/// integer, floats, std::string and bool types.
+template <typename Ty> class Envar {
+  Ty Data;
+  bool IsPresent;
+  bool Initialized;
+
+public:
+  /// Auxiliary function to safely create envars. This static function safely
+  /// creates envars using fallible constructors. See the constructors to know
+  /// more details about the creation parameters.
+  template <typename... ArgsTy>
+  static Expected<Envar> create(ArgsTy &&...Args) {
+    Error Err = Error::success();
+    Envar Envar(std::forward<ArgsTy>(Args)..., Err);
+    if (Err)
+      return std::move(Err);
+    return std::move(Envar);
+  }
+
+  /// Create an empty envar. Cannot be consulted. This constructor is merely
+  /// for convenience. This constructor is not fallible.
+  Envar() : Data(Ty()), IsPresent(false), Initialized(false) {}
+
+  /// Create an envar with a name and an optional default. The Envar object will
+  /// take the value read from the environment variable, or the default if it
+  /// was not set or not correct. This constructor is not fallible.
+  Envar(StringRef Name, Ty Default = Ty())
+      : Data(Default), IsPresent(false), Initialized(true) {
+
+    if (const char *EnvStr = getenv(Name.data())) {
+      // Check whether the envar is defined and valid.
+      IsPresent = StringParser::parse<Ty>(EnvStr, Data);
+
+      if (!IsPresent) {
+        DP("Ignoring invalid value %s for envar %s\n", EnvStr, Name.data());
+        Data = Default;
+      }
+    }
+  }
+
+  /// Get the definitive value.
+  const Ty &get() const {
+    // Throw a runtime error in case this envar is not initialized.
+    if (!Initialized)
+      FATAL_MESSAGE0(1, "Consulting envar before initialization");
+
+    return Data;
+  }
+
+  /// Get the definitive value.
+  operator Ty() const { return get(); }
+
+  /// Indicate whether the environment variable was defined and valid.
+  bool isPresent() const { return IsPresent; }
+
+private:
+  /// This constructor should never fail but we provide it for convenience. This
+  /// way, the constructor can be used by the Envar::create() static function
+  /// to safely create this kind of envars.
+  Envar(StringRef Name, Ty Default, Error &Err) : Envar(Name, Default) {
+    ErrorAsOutParameter EAO(&Err);
+    Err = Error::success();
+  }
+
+  /// Create an envar with a name, getter function and a setter function. The
+  /// Envar object will take the value read from the environment variable if
+  /// this value is accepted by the setter function. Otherwise, the getter
+  /// function will be executed to get the default value. The getter should be
+  /// of the form Error GetterFunctionTy(Ty &Value) and the setter should
+  /// be of the form Error SetterFunctionTy(Ty Value). This constructor has a
+  /// private visibility because is a fallible constructor. Please use the
+  /// Envar::create() static function to safely create this object instead.
+  template <typename GetterFunctor, typename SetterFunctor>
+  Envar(StringRef Name, GetterFunctor Getter, SetterFunctor Setter, Error &Err)
+      : Data(Ty()), IsPresent(false), Initialized(true) {
+    ErrorAsOutParameter EAO(&Err);
+    Err = init(Name, Getter, Setter);
+  }
+
+  template <typename GetterFunctor, typename SetterFunctor>
+  Error init(StringRef Name, GetterFunctor Getter, SetterFunctor Setter);
+};
+
+/// Define some common envar types.
+using IntEnvar = Envar<int>;
+using Int32Envar = Envar<int32_t>;
+using Int64Envar = Envar<int64_t>;
+using UInt32Envar = Envar<uint32_t>;
+using UInt64Envar = Envar<uint64_t>;
+using StringEnvar = Envar<std::string>;
+using BoolEnvar = Envar<bool>;
+
+template <>
+inline bool StringParser::parse(const char *ValueStr, bool &Result) {
+  std::string Value(ValueStr);
+
+  // Convert the string to lowercase.
+  std::transform(Value.begin(), Value.end(), Value.begin(),
+                 [](unsigned char c) { return std::tolower(c); });
+
+  // May be implemented with fancier C++ features, but let's keep it simple.
+  if (Value == "true" || Value == "yes" || Value == "on" || Value == "1")
+    Result = true;
+  else if (Value == "false" || Value == "no" || Value == "off" || Value == "0")
+    Result = false;
+  else
+    return false;
+
+  // Parsed correctly.
+  return true;
+}
+
+template <typename Ty>
+inline bool StringParser::parse(const char *Value, Ty &Result) {
+  assert(Value && "Parsed value cannot be null");
+
+  std::istringstream Stream(Value);
+  Stream >> Result;
+
+  return !Stream.fail();
+}
+
+template <typename Ty>
+template <typename GetterFunctor, typename SetterFunctor>
+inline Error Envar<Ty>::init(StringRef Name, GetterFunctor Getter,
+                             SetterFunctor Setter) {
+  // Get the default value.
+  Ty Default;
+  if (Error Err = Getter(Default))
+    return Err;
+
+  if (const char *EnvStr = getenv(Name.data())) {
+    IsPresent = StringParser::parse<Ty>(EnvStr, Data);
+    if (IsPresent) {
+      // Check whether the envar value is actually valid.
+      Error Err = Setter(Data);
+      if (Err) {
+        // The setter reported an invalid value. Mark the user-defined value as
+        // not present and reset to the getter value (default).
+        IsPresent = false;
+        Data = Default;
+        DP("Setter of envar %s failed, resetting to %s\n", Name.data(),
+           std::to_string(Data).data());
+        consumeError(std::move(Err));
+      }
+    } else {
+      DP("Ignoring invalid value %s for envar %s\n", EnvStr, Name.data());
+      Data = Default;
+    }
+  } else {
+    Data = Default;
+  }
+
+  return Error::success();
+}
+
+} // namespace target
+} // namespace omp
+} // namespace llvm
+
+#endif // OPENMP_LIBOMPTARGET_INCLUDE_UTILITIES_H

openmp/libomptarget/include/rtl.h

@@ -169,6 +169,9 @@ struct RTLsTy {
   // (i.e. the library attempts to load the RTLs (plugins) only once).
   std::once_flag InitFlag;
   void loadRTLs(); // not thread-safe
+
+private:
+  static bool attemptLoadRTL(const std::string &RTLName, RTLInfoTy &RTL);
 };
 
 /// Map between the host entry begin and the translation table. Each

openmp/libomptarget/plugins-nextgen/CMakeLists.txt

@@ -0,0 +1,87 @@
+##===----------------------------------------------------------------------===##
+#
+# 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
+#
+##===----------------------------------------------------------------------===##
+#
+# Build plugins for the user system if available.
+#
+##===----------------------------------------------------------------------===##
+
+add_subdirectory(common)
+
+# void build_generic_elf64_nextgen(string tmachine, string tmachine_name, string tmachine_libname, string elf_machine_id);
+# - build a plugin for an ELF based generic 64-bit target based on libffi.
+# - tmachine: name of the machine processor as used in the cmake build system.
+# - tmachine_name: name of the machine to be printed with the debug messages.
+# - tmachine_libname: machine name to be appended to the plugin library name.
+macro(build_generic_elf64_nextgen tmachine tmachine_name tmachine_libname tmachine_triple elf_machine_id)
+if(CMAKE_SYSTEM_PROCESSOR MATCHES "${tmachine}$")
+  if(LIBOMPTARGET_DEP_LIBFFI_FOUND)
+
+    libomptarget_say("Building ${tmachine_name} NextGen offloading plugin.")
+
+    # Define macro to be used as prefix of the runtime messages for this target.
+    add_definitions("-DTARGET_NAME=${tmachine_name}")
+
+    # Define debug prefix. TODO: This should be automatized in the Debug.h but
+    # it requires changing the original plugins.
+    add_definitions(-DDEBUG_PREFIX="TARGET ${tmachine_name} RTL")
+
+    # Define macro with the ELF ID for this target.
+    add_definitions("-DTARGET_ELF_ID=${elf_machine_id}")
+
+    add_llvm_library("omptarget.rtl.${tmachine_libname}.nextgen"
+      SHARED
+
+      ${CMAKE_CURRENT_SOURCE_DIR}/../generic-elf-64bit/src/rtl.cpp
+
+      ADDITIONAL_HEADER_DIRS
+      ${LIBOMPTARGET_INCLUDE_DIR}
+      ${LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIR}
+
+        LINK_LIBS
+        PRIVATE
+        elf_common
+        MemoryManager
+        PluginInterface
+        ${LIBOMPTARGET_DEP_LIBFFI_LIBRARIES}
+        ${OPENMP_PTHREAD_LIB}
+        "-Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exports"
+
+      NO_INSTALL_RPATH
+    )
+
+    # Install plugin under the lib destination folder.
+    install(TARGETS "omptarget.rtl.${tmachine_libname}.nextgen"
+      LIBRARY DESTINATION "${OPENMP_INSTALL_LIBDIR}")
+    set_target_properties("omptarget.rtl.${tmachine_libname}.nextgen" PROPERTIES
+      INSTALL_RPATH "$ORIGIN" BUILD_RPATH "$ORIGIN:${CMAKE_CURRENT_BINARY_DIR}/.."
+      CXX_VISIBILITY_PRESET protected)
+
+    target_include_directories( "omptarget.rtl.${tmachine_libname}.nextgen" PRIVATE
+      ${LIBOMPTARGET_INCLUDE_DIR}
+      ${LIBOMPTARGET_DEP_LIBFFI_INCLUDE_DIR})
+
+    list(APPEND LIBOMPTARGET_TESTED_PLUGINS
+      "omptarget.rtl.${tmachine_libname}.nextgen")
+
+  else(LIBOMPTARGET_DEP_LIBFFI_FOUND)
+    libomptarget_say("Not building ${tmachine_name} NextGen offloading plugin: libffi dependency not found.")
+  endif(LIBOMPTARGET_DEP_LIBFFI_FOUND)
+else()
+  libomptarget_say("Not building ${tmachine_name} NextGen offloading plugin: machine not found in the system.")
+endif()
+endmacro()
+
+add_subdirectory(aarch64)
+add_subdirectory(cuda)
+add_subdirectory(ppc64)
+add_subdirectory(ppc64le)
+add_subdirectory(x86_64)
+
+# Make sure the parent scope can see the plugins that will be created.
+set(LIBOMPTARGET_SYSTEM_TARGETS "${LIBOMPTARGET_SYSTEM_TARGETS}" PARENT_SCOPE)
+set(LIBOMPTARGET_TESTED_PLUGINS "${LIBOMPTARGET_TESTED_PLUGINS}" PARENT_SCOPE)

openmp/libomptarget/plugins-nextgen/aarch64/CMakeLists.txt

@@ -0,0 +1,17 @@
+##===----------------------------------------------------------------------===##
+#
+# 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
+#
+##===----------------------------------------------------------------------===##
+#
+# Build a plugin for an aarch64 machine if available.
+#
+##===----------------------------------------------------------------------===##
+
+if(CMAKE_SYSTEM_NAME MATCHES "Linux")
+  build_generic_elf64_nextgen("aarch64" "aarch64" "aarch64" "aarch64-unknown-linux-gnu" "183")
+else()
+ libomptarget_say("Not building aarch64 NextGen offloading plugin: machine not found in the system.")
+endif()

openmp/libomptarget/plugins-nextgen/common/CMakeLists.txt

@@ -0,0 +1,13 @@
+##===----------------------------------------------------------------------===##
+#
+# 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
+#
+##===----------------------------------------------------------------------===##
+#
+# Common parts which can be used by all plugins
+#
+##===----------------------------------------------------------------------===##
+
+add_subdirectory(PluginInterface)