DPC++ Runtime
Runtime libraries for oneAPI DPC++
jit_compiler.hpp
Go to the documentation of this file.
1 //==--- jit_compiler.hpp - SYCL runtime JIT compiler for kernel fusion -----==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
14 #include <sycl/feature_test.hpp>
15 #if SYCL_EXT_JIT_ENABLE
16 #include <KernelFusion.h>
17 #endif // SYCL_EXT_JIT_ENABLE
18 
19 #include <unordered_map>
20 
21 namespace jit_compiler {
22 enum class BinaryFormat : uint32_t;
23 class JITContext;
24 struct SYCLKernelInfo;
25 struct SYCLKernelAttribute;
26 template <typename T> class DynArray;
27 using ArgUsageMask = DynArray<uint8_t>;
28 using JITEnvVar = DynArray<char>;
29 } // namespace jit_compiler
30 
31 namespace sycl {
32 inline namespace _V1 {
33 namespace detail {
34 
35 class jit_compiler {
36 
37 public:
38  std::unique_ptr<detail::CG>
39  fuseKernels(QueueImplPtr Queue, std::vector<ExecCGCommand *> &InputKernels,
40  const property_list &);
41  ur_kernel_handle_t
43  const RTDeviceBinaryImage *BinImage,
44  const std::string &KernelName,
45  const std::vector<unsigned char> &SpecConstBlob);
46 
47  bool isAvailable() { return Available; }
48 
50  static jit_compiler instance{};
51  return instance;
52  }
53 
54 private:
55  jit_compiler();
56  ~jit_compiler() = default;
57  jit_compiler(const jit_compiler &) = delete;
58  jit_compiler(jit_compiler &&) = delete;
59  jit_compiler &operator=(const jit_compiler &) = delete;
60  jit_compiler &operator=(const jit_compiler &&) = delete;
61 
63  createPIDeviceBinary(const ::jit_compiler::SYCLKernelInfo &FusedKernelInfo,
64  ::jit_compiler::BinaryFormat Format);
65 
66  std::vector<uint8_t>
67  encodeArgUsageMask(const ::jit_compiler::ArgUsageMask &Mask) const;
68 
69  std::vector<uint8_t> encodeReqdWorkGroupSize(
70  const ::jit_compiler::SYCLKernelAttribute &Attr) const;
71 
72  // Indicate availability of the JIT compiler
73  bool Available;
74 
75  // Manages the lifetime of the UR structs for device binaries.
76  std::vector<DeviceBinariesCollection> JITDeviceBinaries;
77 
78 #if SYCL_EXT_JIT_ENABLE
79  // Handles to the entry points of the lazily loaded JIT library.
80  using FuseKernelsFuncT = decltype(::jit_compiler::fuseKernels) *;
81  using MaterializeSpecConstFuncT =
82  decltype(::jit_compiler::materializeSpecConstants) *;
83  using ResetConfigFuncT = decltype(::jit_compiler::resetJITConfiguration) *;
84  using AddToConfigFuncT = decltype(::jit_compiler::addToJITConfiguration) *;
85  FuseKernelsFuncT FuseKernelsHandle = nullptr;
86  MaterializeSpecConstFuncT MaterializeSpecConstHandle = nullptr;
87  ResetConfigFuncT ResetConfigHandle = nullptr;
88  AddToConfigFuncT AddToConfigHandle = nullptr;
89 #endif // SYCL_EXT_JIT_ENABLE
90 };
91 
92 } // namespace detail
93 } // namespace _V1
94 } // namespace sycl
std::unique_ptr< detail::CG > fuseKernels(QueueImplPtr Queue, std::vector< ExecCGCommand * > &InputKernels, const property_list &)
static jit_compiler & get_instance()
ur_kernel_handle_t materializeSpecConstants(QueueImplPtr Queue, const RTDeviceBinaryImage *BinImage, const std::string &KernelName, const std::vector< unsigned char > &SpecConstBlob)
Objects of the property_list class are containers for the SYCL properties.
DynArray< uint8_t > ArgUsageMask
std::shared_ptr< sycl::detail::queue_impl > QueueImplPtr
Definition: helpers.hpp:45
annotated_arg & operator=(annotated_arg &)=default
Definition: access.hpp:18
This struct is a record of all the device code that may be offloaded.
Definition: compiler.hpp:186