program_manager.hpp

Go to the documentation of this file.

//==------ program_manager.hpp --- SYCL program manager---------------------==//
//
// 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
//
//===----------------------------------------------------------------------===//
 
#pragma once
#include <detail/cg.hpp>
#include <detail/device_binary_image.hpp>
#include <detail/device_global_map_entry.hpp>
#include <detail/host_pipe_map_entry.hpp>
#include <detail/kernel_arg_mask.hpp>
#include <detail/spec_constant_impl.hpp>
#include <sycl/detail/cg_types.hpp>
#include <sycl/detail/common.hpp>
#include <sycl/detail/device_global_map.hpp>
#include <sycl/detail/export.hpp>
#include <sycl/detail/host_pipe_map.hpp>
#include <sycl/detail/os_util.hpp>
#include <sycl/detail/ur.hpp>
#include <sycl/detail/util.hpp>
#include <sycl/device.hpp>
#include <sycl/kernel_bundle.hpp>
 
#include <cstdint>
#include <map>
#include <memory>
#include <set>
#include <unordered_map>
#include <unordered_set>
#include <vector>
 
// +++ Entry points referenced by the offload wrapper object {
 
extern "C" __SYCL_EXPORT void __sycl_register_lib(sycl_device_binaries desc);
 
extern "C" __SYCL_EXPORT void __sycl_unregister_lib(sycl_device_binaries desc);
 
// +++ }
 
namespace sycl {
inline namespace _V1 {
class context;
namespace detail {
 
bool doesDevSupportDeviceRequirements(const device &Dev,
                                      const RTDeviceBinaryImage &BinImages);
std::optional<sycl::exception>
checkDevSupportDeviceRequirements(const device &Dev,
                                  const RTDeviceBinaryImage &BinImages,
                                  const NDRDescT &NDRDesc = {});
 
// This value must be the same as in libdevice/device_itt.h.
// See sycl/doc/design/ITTAnnotations.md for more info.
static constexpr uint32_t inline ITTSpecConstId = 0xFF747469;
 
class context_impl;
using ContextImplPtr = std::shared_ptr<context_impl>;
class device_impl;
using DeviceImplPtr = std::shared_ptr<device_impl>;
class queue_impl;
class event_impl;
// DeviceLibExt is shared between sycl runtime and sycl-post-link tool.
// If any update is made here, need to sync with DeviceLibExt definition
// in llvm/tools/sycl-post-link/sycl-post-link.cpp
enum class DeviceLibExt : std::uint32_t {
  cl_intel_devicelib_assert,
  cl_intel_devicelib_math,
  cl_intel_devicelib_math_fp64,
  cl_intel_devicelib_complex,
  cl_intel_devicelib_complex_fp64,
  cl_intel_devicelib_cstring,
  cl_intel_devicelib_imf,
  cl_intel_devicelib_imf_fp64,
  cl_intel_devicelib_imf_bf16,
  cl_intel_devicelib_bfloat16,
};
 
// Provides single loading and building OpenCL programs with unique contexts
// that is necessary for no interoperability cases with lambda.
class ProgramManager {
public:
  // Returns the single instance of the program manager for the entire
  // process. Can only be called after staticInit is done.
  static ProgramManager &getInstance();
 
  RTDeviceBinaryImage &getDeviceImage(const std::string &KernelName,
                                      const context &Context,
                                      const device &Device,
                                      bool JITCompilationIsRequired = false);
 
  RTDeviceBinaryImage &getDeviceImage(
      const std::unordered_set<RTDeviceBinaryImage *> &ImagesToVerify,
      const context &Context, const device &Device,
      bool JITCompilationIsRequired = false);
 
  ur_program_handle_t createURProgram(const RTDeviceBinaryImage &Img,
                                      const context &Context,
                                      const device &Device);
  std::pair<ur_program_handle_t, bool>
  getOrCreateURProgram(
      const RTDeviceBinaryImage &Img,
      const std::vector<const RTDeviceBinaryImage *> &AllImages,
      const context &Context,
      const device &Device,
      const std::string &CompileAndLinkOptions,
      SerializedObj SpecConsts);
  ur_program_handle_t getBuiltURProgram(const ContextImplPtr &ContextImpl,
                                        const DeviceImplPtr &DeviceImpl,
                                        const std::string &KernelName,
                                        const NDRDescT &NDRDesc = {},
                                        bool JITCompilationIsRequired = false);
 
  ur_program_handle_t getBuiltURProgram(const context &Context,
                                        const device &Device,
                                        const std::string &KernelName,
                                        const property_list &PropList,
                                        bool JITCompilationIsRequired = false);
 
  std::tuple<ur_kernel_handle_t, std::mutex *, const KernelArgMask *,
             ur_program_handle_t>
  getOrCreateKernel(const ContextImplPtr &ContextImpl,
                    const DeviceImplPtr &DeviceImpl,
                    const std::string &KernelName,
                    const NDRDescT &NDRDesc = {});
 
  ur_kernel_handle_t getCachedMaterializedKernel(
      const std::string &KernelName,
      const std::vector<unsigned char> &SpecializationConsts);
 
  ur_kernel_handle_t getOrCreateMaterializedKernel(
      const RTDeviceBinaryImage &Img, const context &Context,
      const device &Device, const std::string &KernelName,
      const std::vector<unsigned char> &SpecializationConsts);
 
  ur_program_handle_t getUrProgramFromUrKernel(ur_kernel_handle_t Kernel,
                                               const ContextImplPtr Context);
 
  void addImages(sycl_device_binaries DeviceImages);
  void debugPrintBinaryImages() const;
  static std::string getProgramBuildLog(const ur_program_handle_t &Program,
                                        const ContextImplPtr Context);
 
  uint32_t getDeviceLibReqMask(const RTDeviceBinaryImage &Img);
 
  const KernelArgMask *
  getEliminatedKernelArgMask(ur_program_handle_t NativePrg,
                             const std::string &KernelName);
 
  // The function returns the unique SYCL kernel identifier associated with a
  // kernel name.
  kernel_id getSYCLKernelID(const std::string &KernelName);
 
  // The function returns a vector containing all unique SYCL kernel identifiers
  // in SYCL device images.
  std::vector<kernel_id> getAllSYCLKernelIDs();
 
  // The function returns the unique SYCL kernel identifier associated with a
  // built-in kernel name.
  kernel_id getBuiltInKernelID(const std::string &KernelName);
 
  // The function inserts or initializes a device_global entry into the
  // device_global map.
  void addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
                                  const char *UniqueId);
 
  // Returns true if any available image is compatible with the device Dev.
  bool hasCompatibleImage(const device &Dev);
 
  // The function gets a device_global entry identified by the pointer to the
  // device_global object from the device_global map.
  DeviceGlobalMapEntry *getDeviceGlobalEntry(const void *DeviceGlobalPtr);
 
  // The function gets multiple device_global entries identified by their unique
  // IDs from the device_global map.
  std::vector<DeviceGlobalMapEntry *>
  getDeviceGlobalEntries(const std::vector<std::string> &UniqueIds,
                         bool ExcludeDeviceImageScopeDecorated = false);
  // The function inserts or initializes a host_pipe entry into the
  // host_pipe map.
  void addOrInitHostPipeEntry(const void *HostPipePtr, const char *UniqueId);
 
  // The function gets a host_pipe entry identified by the unique ID from
  // the host_pipe map.
  HostPipeMapEntry *getHostPipeEntry(const std::string &UniqueId);
 
  // The function gets a host_pipe entry identified by the pointer to the
  // host_pipe object from the host_pipe map.
  HostPipeMapEntry *getHostPipeEntry(const void *HostPipePtr);
 
  device_image_plain
  getDeviceImageFromBinaryImage(RTDeviceBinaryImage *BinImage,
                                const context &Ctx, const device &Dev);
 
  // The function returns a vector of SYCL device images that are compiled with
  // the required state and at least one device from the passed list of devices.
  std::vector<device_image_plain> getSYCLDeviceImagesWithCompatibleState(
      const context &Ctx, const std::vector<device> &Devs,
      bundle_state TargetState, const std::vector<kernel_id> &KernelIDs = {});
 
  // Brind images in the passed vector to the required state. Does it inplace
  void
  bringSYCLDeviceImagesToState(std::vector<device_image_plain> &DeviceImages,
                               bundle_state TargetState);
 
  // The function returns a vector of SYCL device images in required state,
  // which are compatible with at least one of the device from Devs.
  std::vector<device_image_plain>
  getSYCLDeviceImages(const context &Ctx, const std::vector<device> &Devs,
                      bundle_state State);
 
  // The function returns a vector of SYCL device images, for which Selector
  // callable returns true, in required state, which are compatible with at
  // least one of the device from Devs.
  std::vector<device_image_plain>
  getSYCLDeviceImages(const context &Ctx, const std::vector<device> &Devs,
                      const DevImgSelectorImpl &Selector,
                      bundle_state TargetState);
 
  // The function returns a vector of SYCL device images which represent at
  // least one kernel from kernel ids vector in required state, which are
  // compatible with at least one of the device from Devs.
  std::vector<device_image_plain>
  getSYCLDeviceImages(const context &Ctx, const std::vector<device> &Devs,
                      const std::vector<kernel_id> &KernelIDs,
                      bundle_state TargetState);
 
  // Produces new device image by convering input device image to the object
  // state
  device_image_plain compile(const device_image_plain &DeviceImage,
                             const std::vector<device> &Devs,
                             const property_list &PropList);
 
  // Produces set of device images by convering input device images to object
  // the executable state
  std::vector<device_image_plain> link(const device_image_plain &DeviceImages,
                                       const std::vector<device> &Devs,
                                       const property_list &PropList);
 
  // Produces new device image by converting input device image to the
  // executable state
  device_image_plain build(const device_image_plain &DeviceImage,
                           const std::vector<device> &Devs,
                           const property_list &PropList);
 
  std::tuple<ur_kernel_handle_t, std::mutex *, const KernelArgMask *>
  getOrCreateKernel(const context &Context, const std::string &KernelName,
                    const property_list &PropList, ur_program_handle_t Program);
 
  ProgramManager();
  ~ProgramManager() = default;
 
  bool kernelUsesAssert(const std::string &KernelName) const;
 
  bool kernelUsesAsan() const { return m_AsanFoundInImage; }
 
  std::set<RTDeviceBinaryImage *>
  getRawDeviceImages(const std::vector<kernel_id> &KernelIDs);
 
private:
  ProgramManager(ProgramManager const &) = delete;
  ProgramManager &operator=(ProgramManager const &) = delete;
 
  using ProgramPtr = std::unique_ptr<remove_pointer_t<ur_program_handle_t>,
                                     decltype(&::urProgramRelease)>;
  ProgramPtr build(ProgramPtr Program, const ContextImplPtr Context,
                   const std::string &CompileOptions,
                   const std::string &LinkOptions, ur_device_handle_t Device,
                   uint32_t DeviceLibReqMask,
                   const std::vector<ur_program_handle_t> &ProgramsToLink);
 
  void dumpImage(const RTDeviceBinaryImage &Img, uint32_t SequenceID = 0) const;
 
  void cacheKernelUsesAssertInfo(RTDeviceBinaryImage &Img);
 
  std::set<RTDeviceBinaryImage *>
  collectDeviceImageDepsForImportedSymbols(const RTDeviceBinaryImage &Img,
                                           device Dev);
 
  std::set<RTDeviceBinaryImage *>
  collectDependentDeviceImagesForVirtualFunctions(
      const RTDeviceBinaryImage &Img, device Dev);
 
  using RTDeviceBinaryImageUPtr = std::unique_ptr<RTDeviceBinaryImage>;
 
  //
  std::unordered_map<std::string, kernel_id> m_KernelName2KernelIDs;
 
  // Maps KernelIDs to device binary images. There can be more than one image
  // in case of SPIRV + AOT.
  // Using shared_ptr to avoid expensive copy of the vector.
  std::unordered_multimap<kernel_id, RTDeviceBinaryImage *>
      m_KernelIDs2BinImage;
 
  // Maps device binary image to a vector of kernel ids in this image.
  // Using shared_ptr to avoid expensive copy of the vector.
  // The vector is initialized in addImages function and is supposed to be
  // immutable afterwards.
  std::unordered_map<RTDeviceBinaryImage *,
                     std::shared_ptr<std::vector<kernel_id>>>
      m_BinImg2KernelIDs;
 
  std::mutex m_KernelIDsMutex;
 
  std::unordered_multimap<std::string, RTDeviceBinaryImage *> m_ServiceKernels;
 
  // from kernel bundles.
  std::unordered_multimap<std::string, RTDeviceBinaryImage *>
      m_ExportedSymbolImages;
 
  std::unordered_set<RTDeviceBinaryImageUPtr> m_DeviceImages;
 
  std::unordered_map<std::string, kernel_id> m_BuiltInKernelIDs;
 
  std::unordered_map<std::string, std::set<RTDeviceBinaryImage *>>
      m_VFSet2BinImage;
 
  std::mutex m_BuiltInKernelIDsMutex;
 
  // Keeps track of ur_program to image correspondence. Needed for:
  // - knowing which specialization constants are used in the program and
  //   injecting their current values before compiling the SPIR-V; the binary
  //   image object has info about all spec constants used in the module
  // - finding kernel argument masks for kernels associated with each
  //   ur_program
  // NOTE: using RTDeviceBinaryImage raw pointers is OK, since they are not
  // referenced from outside SYCL runtime and RTDeviceBinaryImage object
  // lifetime matches program manager's one.
  // NOTE: keys in the map can be invalid (reference count went to zero and
  // the underlying program disposed of), so the map can't be used in any way
  // other than binary image lookup with known live UrProgram as the key.
  // NOTE: access is synchronized via the MNativeProgramsMutex
  std::unordered_map<ur_program_handle_t, const RTDeviceBinaryImage *>
      NativePrograms;
 
  std::mutex MNativeProgramsMutex;
 
  using KernelNameToArgMaskMap = std::unordered_map<std::string, KernelArgMask>;
  std::unordered_map<const RTDeviceBinaryImage *, KernelNameToArgMaskMap>
      m_EliminatedKernelArgMasks;
 
  bool m_UseSpvFile = false;
  RTDeviceBinaryImageUPtr m_SpvFileImage;
 
  std::set<std::string> m_KernelUsesAssert;
 
  // True iff there is a device image compiled with AddressSanitizer
  bool m_AsanFoundInImage;
 
  // Maps between device_global identifiers and associated information.
  std::unordered_map<std::string, std::unique_ptr<DeviceGlobalMapEntry>>
      m_DeviceGlobals;
  std::unordered_map<const void *, DeviceGlobalMapEntry *> m_Ptr2DeviceGlobal;
 
  std::mutex m_DeviceGlobalsMutex;
 
  // Maps between host_pipe identifiers and associated information.
  std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>
      m_HostPipes;
  std::unordered_map<const void *, HostPipeMapEntry *> m_Ptr2HostPipe;
 
  std::mutex m_HostPipesMutex;
 
  using MaterializedEntries =
      std::map<std::vector<unsigned char>, ur_kernel_handle_t>;
  std::unordered_map<std::string, MaterializedEntries> m_MaterializedKernels;
};
} // namespace detail
} // namespace _V1
} // namespace sycl