DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_bundle.cpp
Go to the documentation of this file.
1 //==------- kernel_bundle.cpp - SYCL kernel_bundle and free functions ------==//
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 
14 
15 #include <cstddef>
16 #include <set>
17 #include <vector>
18 
19 namespace sycl {
20 inline namespace _V1 {
21 
22 kernel_id::kernel_id(const char *Name)
23  : impl(std::make_shared<detail::kernel_id_impl>(Name)) {}
24 
25 const char *kernel_id::get_name() const noexcept { return impl->get_name(); }
26 
27 namespace detail {
28 
32 
33 bool device_image_plain::has_kernel(const kernel_id &KernelID) const noexcept {
34  return impl->has_kernel(KernelID);
35 }
36 
37 bool device_image_plain::has_kernel(const kernel_id &KernelID,
38  const device &Dev) const noexcept {
39  return impl->has_kernel(KernelID, Dev);
40 }
41 
42 pi_native_handle device_image_plain::getNative() const {
43  return impl->getNative();
44 }
45 
49 
50 bool kernel_bundle_plain::empty() const noexcept { return impl->empty(); }
51 
52 backend kernel_bundle_plain::get_backend() const noexcept {
53  return impl->get_backend();
54 }
55 
56 context kernel_bundle_plain::get_context() const noexcept {
57  return impl->get_context();
58 }
59 
60 std::vector<device> kernel_bundle_plain::get_devices() const noexcept {
61  return impl->get_devices();
62 }
63 
64 std::vector<kernel_id> kernel_bundle_plain::get_kernel_ids() const {
65  return impl->get_kernel_ids();
66 }
67 
68 bool kernel_bundle_plain::contains_specialization_constants() const noexcept {
69  return impl->contains_specialization_constants();
70 }
71 
72 bool kernel_bundle_plain::native_specialization_constant() const noexcept {
73  return impl->native_specialization_constant();
74 }
75 
76 kernel kernel_bundle_plain::get_kernel(const kernel_id &KernelID) const {
77  return impl->get_kernel(KernelID, impl);
78 }
79 
80 const device_image_plain *kernel_bundle_plain::begin() const {
81  return impl->begin();
82 }
83 
84 const device_image_plain *kernel_bundle_plain::end() const {
85  return impl->end();
86 }
87 
88 bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID) const noexcept {
89  return impl->has_kernel(KernelID);
90 }
91 
92 bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID,
93  const device &Dev) const noexcept {
94  return impl->has_kernel(KernelID, Dev);
95 }
96 
97 bool kernel_bundle_plain::has_specialization_constant_impl(
98  const char *SpecName) const noexcept {
99  return impl->has_specialization_constant(SpecName);
100 }
101 
102 void kernel_bundle_plain::set_specialization_constant_impl(
103  const char *SpecName, void *Value, size_t Size) noexcept {
104  impl->set_specialization_constant_raw_value(SpecName, Value, Size);
105 }
106 
107 void kernel_bundle_plain::get_specialization_constant_impl(
108  const char *SpecName, void *Value) const noexcept {
109  impl->get_specialization_constant_raw_value(SpecName, Value);
110 }
111 
112 bool kernel_bundle_plain::is_specialization_constant_set(
113  const char *SpecName) const noexcept {
114  return impl->is_specialization_constant_set(SpecName);
115 }
116 
117 bool kernel_bundle_plain::ext_oneapi_has_kernel(const std::string &name) {
118  return impl->ext_oneapi_has_kernel(name);
119 }
120 
121 kernel kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) {
122  return impl->ext_oneapi_get_kernel(name, impl);
123 }
124 
128 
129 const std::vector<device>
130 removeDuplicateDevices(const std::vector<device> &Devs) {
131  std::vector<device> UniqueDevices;
132 
133  // Building a new vector with unique elements and keep original order
134  std::unordered_set<device> UniqueDeviceSet;
135  for (const device &Dev : Devs)
136  if (UniqueDeviceSet.insert(Dev).second)
137  UniqueDevices.push_back(Dev);
138 
139  return UniqueDevices;
140 }
141 
143 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
144  return detail::ProgramManager::getInstance().getSYCLKernelID(
145  KernelName.data());
146 #else
147  return detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
148 #endif
149 }
150 
152 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
153  bundle_state State) {
154  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, State);
155 }
156 
158 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
159  const std::vector<kernel_id> &KernelIDs,
160  bundle_state State) {
161  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, KernelIDs,
162  State);
163 }
164 
166 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
167  bundle_state State, const DevImgSelectorImpl &Selector) {
168  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, Selector,
169  State);
170 }
171 
174  const std::vector<device> &Devs) {
175  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
176 }
177 
178 std::shared_ptr<detail::kernel_bundle_impl>
179 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
180  bundle_state State) {
181  return std::make_shared<detail::kernel_bundle_impl>(Bundles, State);
182 }
183 
184 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
185  bundle_state State) {
186  // Check that all requested devices are associated with the context
187  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
188  if (Devs.empty() || !AllDevicesInTheContext)
189  throw sycl::exception(make_error_code(errc::invalid),
190  "Not all devices are associated with the context or "
191  "vector of devices is empty");
192 
193  if (bundle_state::input == State &&
194  !checkAllDevicesHaveAspect(Devs, aspect::online_compiler))
195  return false;
196  if (bundle_state::object == State &&
197  !checkAllDevicesHaveAspect(Devs, aspect::online_linker))
198  return false;
199 
200  const std::vector<device_image_plain> DeviceImages =
201  detail::ProgramManager::getInstance()
202  .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
203 
204  return (bool)DeviceImages.size();
205 }
206 
207 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
208  const std::vector<kernel_id> &KernelIds,
209  bundle_state State) {
210  // Check that all requested devices are associated with the context
211  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
212 
213  if (Devs.empty() || !AllDevicesInTheContext)
214  throw sycl::exception(make_error_code(errc::invalid),
215  "Not all devices are associated with the context or "
216  "vector of devices is empty");
217 
218  bool DeviceHasRequireAspectForState = true;
219  if (bundle_state::input == State) {
220  DeviceHasRequireAspectForState =
221  std::all_of(Devs.begin(), Devs.end(), [](const device &Dev) {
222  return Dev.has(aspect::online_compiler);
223  });
224  } else if (bundle_state::object == State) {
225  DeviceHasRequireAspectForState =
226  std::all_of(Devs.begin(), Devs.end(), [](const device &Dev) {
227  return Dev.has(aspect::online_linker);
228  });
229  }
230 
231  if (!DeviceHasRequireAspectForState)
232  return false;
233 
234  const std::vector<device_image_plain> DeviceImages =
235  detail::ProgramManager::getInstance()
236  .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
237 
238  std::set<kernel_id, LessByNameComp> CombinedKernelIDs;
239  for (const device_image_plain &DeviceImage : DeviceImages) {
240  const std::shared_ptr<device_image_impl> &DeviceImageImpl =
241  getSyclObjImpl(DeviceImage);
242 
243  CombinedKernelIDs.insert(DeviceImageImpl->get_kernel_ids_ptr()->begin(),
244  DeviceImageImpl->get_kernel_ids_ptr()->end());
245  }
246 
247  const bool AllKernelIDsRepresented =
248  std::all_of(KernelIds.begin(), KernelIds.end(),
249  [&CombinedKernelIDs](const kernel_id &KernelID) {
250  return CombinedKernelIDs.count(KernelID);
251  });
252 
253  return AllKernelIDsRepresented;
254 }
255 
256 std::shared_ptr<detail::kernel_bundle_impl>
258  const std::vector<device> &Devs, const property_list &PropList) {
259  return std::make_shared<detail::kernel_bundle_impl>(
260  InputBundle, Devs, PropList, bundle_state::object);
261 }
262 
263 std::shared_ptr<detail::kernel_bundle_impl>
264 link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
265  const std::vector<device> &Devs, const property_list &PropList) {
266  return std::make_shared<detail::kernel_bundle_impl>(ObjectBundles, Devs,
267  PropList);
268 }
269 
270 std::shared_ptr<detail::kernel_bundle_impl>
272  const std::vector<device> &Devs, const property_list &PropList) {
273  return std::make_shared<detail::kernel_bundle_impl>(
274  InputBundle, Devs, PropList, bundle_state::executable);
275 }
276 
277 // This function finds intersection of associated devices in common for all
278 // bundles
279 std::vector<sycl::device> find_device_intersection(
280  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles) {
281  std::vector<sycl::device> IntersectDevices;
282  std::vector<unsigned int> DevsCounters;
283  std::map<device, unsigned int, LessByHash<device>> DevCounters;
284  for (const sycl::kernel_bundle<bundle_state::object> &ObjectBundle :
285  ObjectBundles)
286  // Increment counter in "DevCounters" each time a device is seen
287  for (const sycl::device &Device : ObjectBundle.get_devices())
288  DevCounters[Device]++;
289 
290  // If some device counter is less than ObjectBundles.size() then some bundle
291  // doesn't have it - do not add such a device to the final result
292  for (const std::pair<const device, unsigned int> &It : DevCounters)
293  if (ObjectBundles.size() == It.second)
294  IntersectDevices.push_back(It.first);
295 
296  return IntersectDevices;
297 }
298 
299 } // namespace detail
300 
304 
305 std::vector<kernel_id> get_kernel_ids() {
306  return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
307 }
308 
309 bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
310  if (KernelIDs.empty())
311  return true;
312  // TODO: also need to check that the architecture specified by the
313  // "-fsycl-targets" flag matches the device when we are able to get the
314  // device's arch.
315  auto doesImageTargetMatchDevice = [](const device &Dev,
316  const detail::RTDeviceBinaryImage &Img) {
317  const char *Target = Img.getRawData().DeviceTargetSpec;
318  auto BE = Dev.get_backend();
319  // ESIMD emulator backend is only compatible with esimd kernels.
320  if (BE == sycl::backend::ext_intel_esimd_emulator) {
321  pi_device_binary_property Prop = Img.getProperty("isEsimdImage");
322  return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0));
323  }
324  if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) {
325  return (BE == sycl::backend::opencl ||
326  BE == sycl::backend::ext_oneapi_level_zero);
327  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) ==
328  0) {
329  return Dev.is_cpu();
330  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN) ==
331  0) {
332  return Dev.is_gpu() && (BE == sycl::backend::opencl ||
333  BE == sycl::backend::ext_oneapi_level_zero);
334  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA) ==
335  0) {
336  return Dev.is_accelerator();
337  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64) == 0) {
338  return BE == sycl::backend::ext_oneapi_cuda;
339  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN) == 0) {
340  return BE == sycl::backend::ext_oneapi_hip;
341  }
342 
343  return false;
344  };
345 
346  // One kernel may be contained in several binary images depending on the
347  // number of targets. This kernel is compatible with the device if there is
348  // at least one image (containing this kernel) whose aspects are supported by
349  // the device and whose target matches the device.
350  for (const auto &KernelID : KernelIDs) {
351  std::set<detail::RTDeviceBinaryImage *> BinImages =
352  detail::ProgramManager::getInstance().getRawDeviceImages({KernelID});
353 
354  if (std::none_of(BinImages.begin(), BinImages.end(),
355  [&](const detail::RTDeviceBinaryImage *Img) {
356  return doesDevSupportDeviceRequirements(Dev, *Img) &&
357  doesImageTargetMatchDevice(Dev, *Img);
358  }))
359  return false;
360  }
361 
362  return true;
363 }
364 
366 // * kernel_compiler extension *
368 namespace ext::oneapi::experimental {
369 
373 
375 // syclex::is_source_kernel_bundle_supported
378  // Support is limited to the opencl and level_zero backends.
379  bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
380  (BE == sycl::backend::opencl);
381  if (BE_Acceptable) {
382  // At the moment, OpenCL and SPIR-V are the only supported languages.
383  if (Language == source_language::opencl) {
385  } else if (Language == source_language::spirv) {
386  return true;
387  }
388  }
389 
390  // otherwise
391  return false;
392 }
393 
395 // syclex::create_kernel_bundle_from_source
397 
399  source_language Language,
400  const std::string &Source) {
401  // TODO: if we later support a "reason" why support isn't present
402  // (like a missing shared library etc.) it'd be nice to include it in
403  // the exception message here.
404  backend BE = SyclContext.get_backend();
405  if (!is_source_kernel_bundle_supported(BE, Language))
406  throw sycl::exception(make_error_code(errc::invalid),
407  "kernel_bundle creation from source not supported");
408 
409  std::shared_ptr<kernel_bundle_impl> KBImpl =
410  std::make_shared<kernel_bundle_impl>(SyclContext, Language, Source);
411  return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
412 }
413 
414 source_kb
416  source_language Language,
417  const std::vector<std::byte> &Bytes) {
418  backend BE = SyclContext.get_backend();
419  if (!is_source_kernel_bundle_supported(BE, Language))
420  throw sycl::exception(make_error_code(errc::invalid),
421  "kernel_bundle creation from source not supported");
422 
423  std::shared_ptr<kernel_bundle_impl> KBImpl =
424  std::make_shared<kernel_bundle_impl>(SyclContext, Language, Bytes);
425  return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
426 }
427 
429 // syclex::detail::build_from_source(source_kb) => exe_kb
431 namespace detail {
432 
434  const std::vector<device> &Devices,
435  const std::vector<std::string> &BuildOptions,
436  std::string *LogPtr) {
437  std::vector<device> UniqueDevices =
439  std::shared_ptr<kernel_bundle_impl> sourceImpl = getSyclObjImpl(SourceKB);
440  std::shared_ptr<kernel_bundle_impl> KBImpl =
441  sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr);
442  return sycl::detail::createSyclObjFromImpl<exe_kb>(KBImpl);
443 }
444 
445 } // namespace detail
446 } // namespace ext::oneapi::experimental
447 
448 } // namespace _V1
449 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:51
backend get_backend() const noexcept
Returns the backend associated with this context.
Definition: context.cpp:132
const char * data() const noexcept
Definition: string_view.hpp:38
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:66
bool is_accelerator() const
Check if device is an accelerator device.
Definition: device.cpp:91
bool is_gpu() const
Check if device is a GPU device.
Definition: device.cpp:89
static std::vector< device > get_devices(info::device_type deviceType=info::device_type::all)
Query available SYCL devices.
Definition: device.cpp:51
backend get_backend() const noexcept
Returns the backend associated with this device.
Definition: device.cpp:237
bool is_cpu() const
Check if device is a CPU device.
Definition: device.cpp:87
The kernel_bundle class represents collection of device images in a particular state.
Objects of the class identify kernel is some kernel_bundle related APIs.
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:79
Objects of the property_list class are containers for the SYCL properties.
std::vector< sycl::device > find_device_intersection(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles)
std::shared_ptr< detail::kernel_bundle_impl > build_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
const std::vector< device > removeDuplicateDevices(const std::vector< device > &Devs)
detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
std::shared_ptr< detail::kernel_bundle_impl > compile_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
std::shared_ptr< detail::kernel_bundle_impl > join_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
kernel_id get_kernel_id_impl(string_view KernelName)
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
std::string string
Definition: handler.hpp:423
bool has_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
detail::KernelBundleImplPtr get_empty_interop_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs)
std::shared_ptr< detail::kernel_bundle_impl > link_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, const std::vector< device > &Devs, const property_list &PropList)
exe_kb build_from_source(source_kb &SourceKB, const std::vector< device > &Devices, const std::vector< std::string > &BuildOptions, std::string *LogPtr)
bool is_source_kernel_bundle_supported(backend BE, source_language Language)
kernel_bundle< bundle_state::ext_oneapi_source > create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::string &Source)
sycl::detail::kernel_bundle_impl kernel_bundle_impl
kernel_bundle< sycl::bundle_state::ext_oneapi_source > source_kb
std::vector< kernel_id > get_kernel_ids()
bool is_compatible(const std::vector< kernel_id > &KernelIDs, const device &Dev)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:107
Definition: access.hpp:18
uintptr_t pi_native_handle
Definition: pi.h:206
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64
Device-specific binary images produced from SPIR-V 64-bit <-> various "spir64_*" triples for specific...
Definition: pi.h:942
#define __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64
PTX 64-bit image <-> "nvptx64", 64-bit NVIDIA PTX device.
Definition: pi.h:946
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64
SPIR-V 64-bit image <-> "spir64", 64-bit OpenCL device.
Definition: pi.h:939
#define __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN
Definition: pi.h:947
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN
Definition: pi.h:943
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA
Definition: pi.h:944
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept