DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_bundle.hpp
Go to the documentation of this file.
1 //==------- kernel_bundle.hpp - 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 
9 #pragma once
10 
11 #include <sycl/context.hpp>
12 #include <sycl/detail/common.hpp>
14 #include <sycl/detail/pi.h>
15 #include <sycl/detail/pi.hpp>
16 #include <sycl/device.hpp>
17 #include <sycl/kernel.hpp>
19 
20 #include <cassert>
21 #include <memory>
22 #include <set>
23 #include <vector>
24 
25 namespace sycl {
27 // Forward declaration
28 template <backend Backend> class backend_traits;
29 template <backend Backend, bundle_state State>
30 auto get_native(const kernel_bundle<State> &Obj)
31  -> backend_return_t<Backend, kernel_bundle<State>>;
32 
33 namespace detail {
34 class kernel_id_impl;
35 }
36 
40 class __SYCL_EXPORT kernel_id {
41 public:
42  kernel_id() = delete;
43 
45  const char *get_name() const noexcept;
46 
47  bool operator==(const kernel_id &RHS) const { return impl == RHS.impl; }
48 
49  bool operator!=(const kernel_id &RHS) const { return !(*this == RHS); }
50 
51 private:
52  kernel_id(const char *Name);
53 
54  kernel_id(const std::shared_ptr<detail::kernel_id_impl> &Impl)
55  : impl(std::move(Impl)) {}
56 
57  std::shared_ptr<detail::kernel_id_impl> impl;
58 
59  template <class Obj>
60  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
61 
62  template <class T>
63  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
64 };
65 
66 namespace detail {
67 class device_image_impl;
68 using DeviceImageImplPtr = std::shared_ptr<device_image_impl>;
69 
70 // The class is used as a base for device_image for "untemplating" public
71 // methods.
72 class __SYCL_EXPORT device_image_plain {
73 public:
75  : impl(std::move(Impl)) {}
76 
77  bool operator==(const device_image_plain &RHS) const {
78  return impl == RHS.impl;
79  }
80 
81  bool operator!=(const device_image_plain &RHS) const {
82  return !(*this == RHS);
83  }
84 
85  bool has_kernel(const kernel_id &KernelID) const noexcept;
86 
87  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
88 
89  pi_native_handle getNative() const;
90 
91 protected:
93 
94  template <class Obj>
95  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
96 
97  template <class T>
98  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
99 };
100 } // namespace detail
101 
103 template <sycl::bundle_state State>
105 public:
106  device_image() = delete;
107 
110  bool has_kernel(const kernel_id &KernelID) const noexcept {
111  return device_image_plain::has_kernel(KernelID);
112  }
113 
116  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
117  return device_image_plain::has_kernel(KernelID, Dev);
118  }
119 
120 private:
122  : device_image_plain(std::move(Impl)) {}
123 
124  template <class Obj>
125  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
126 
127  template <class T>
128  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
129 };
130 
131 namespace detail {
132 class kernel_bundle_impl;
133 using KernelBundleImplPtr = std::shared_ptr<detail::kernel_bundle_impl>;
134 
135 // The class is used as a base for kernel_bundle to "untemplate" it's methods
136 class __SYCL_EXPORT kernel_bundle_plain {
137 public:
139  : impl(std::move(Impl)) {}
140 
141  bool operator==(const kernel_bundle_plain &RHS) const {
142  return impl == RHS.impl;
143  }
144 
145  bool operator!=(const kernel_bundle_plain &RHS) const {
146  return !(*this == RHS);
147  }
148 
149  bool empty() const noexcept;
150 
151  backend get_backend() const noexcept;
152 
153  context get_context() const noexcept;
154 
155  std::vector<device> get_devices() const noexcept;
156 
157  bool has_kernel(const kernel_id &KernelID) const noexcept;
158 
159  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
160 
161  std::vector<kernel_id> get_kernel_ids() const;
162 
163  bool contains_specialization_constants() const noexcept;
164 
165  bool native_specialization_constant() const noexcept;
166 
167 protected:
168  // \returns a kernel object which represents the kernel identified by
169  // kernel_id passed
170  kernel get_kernel(const kernel_id &KernelID) const;
171 
172  // \returns an iterator to the first device image kernel_bundle contains
173  const device_image_plain *begin() const;
174 
175  // \returns an iterator to the last device image kernel_bundle contains
176  const device_image_plain *end() const;
177 
178  bool has_specialization_constant_impl(const char *SpecName) const noexcept;
179 
180  void set_specialization_constant_impl(const char *SpecName, void *Value,
181  size_t Size) noexcept;
182 
183  void get_specialization_constant_impl(const char *SpecName,
184  void *Value) const noexcept;
185 
186  bool is_specialization_constant_set(const char *SpecName) const noexcept;
187 
188  detail::KernelBundleImplPtr impl;
189 };
190 
191 } // namespace detail
192 
197 template <bundle_state State>
198 class kernel_bundle : public detail::kernel_bundle_plain {
199 public:
201 
202  kernel_bundle() = delete;
203 
205  bool empty() const noexcept { return kernel_bundle_plain::empty(); }
206 
208  backend get_backend() const noexcept {
209  return kernel_bundle_plain::get_backend();
210  }
211 
213  context get_context() const noexcept {
214  return kernel_bundle_plain::get_context();
215  }
216 
218  std::vector<device> get_devices() const noexcept {
219  return kernel_bundle_plain::get_devices();
220  }
221 
224  bool has_kernel(const kernel_id &KernelID) const noexcept {
225  return kernel_bundle_plain::has_kernel(KernelID);
226  }
227 
231  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
232  return kernel_bundle_plain::has_kernel(KernelID, Dev);
233  }
234 
236  std::vector<kernel_id> get_kernel_ids() const {
238  }
239 
242  bool contains_specialization_constants() const noexcept {
243  return kernel_bundle_plain::contains_specialization_constants();
244  }
245 
248  bool native_specialization_constant() const noexcept {
249  return kernel_bundle_plain::native_specialization_constant();
250  }
251 
254  template <bundle_state _State = State,
256  kernel get_kernel(const kernel_id &KernelID) const {
257  return detail::kernel_bundle_plain::get_kernel(KernelID);
258  }
259 
260  // This guard is needed because the libsycl.so can compiled with C++ <=14
261  // while the code requires C++17. This code is not supposed to be used by the
262  // libsycl.so so it should not be a problem.
263 #if __cplusplus >= 201703L
266  template <auto &SpecName> bool has_specialization_constant() const noexcept {
267  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
268  return has_specialization_constant_impl(SpecSymName);
269  }
270 
274  template <auto &SpecName, bundle_state _State = State,
275  typename = detail::enable_if_t<_State == bundle_state::input>>
276  void set_specialization_constant(
277  typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
278  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
279  set_specialization_constant_impl(SpecSymName, &Value,
280  sizeof(decltype(Value)));
281  }
282 
285  template <auto &SpecName>
286  typename std::remove_reference_t<decltype(SpecName)>::value_type
287  get_specialization_constant() const {
288  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
289  if (!is_specialization_constant_set(SpecSymName))
290  return SpecName.getDefaultValue();
291 
292  using SCType =
293  typename std::remove_reference_t<decltype(SpecName)>::value_type;
294 
295  std::array<char *, sizeof(SCType)> RetValue;
296 
297  get_specialization_constant_impl(SpecSymName, RetValue.data());
298 
299  return *reinterpret_cast<SCType *>(RetValue.data());
300  }
301 #endif
302 
305  return reinterpret_cast<device_image_iterator>(
306  kernel_bundle_plain::begin());
307  }
308 
311  return reinterpret_cast<device_image_iterator>(kernel_bundle_plain::end());
312  }
313 
314 private:
316  : kernel_bundle_plain(std::move(Impl)) {}
317 
318  template <class Obj>
319  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
320 
321  template <class T>
322  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
323 
324  template <backend Backend, bundle_state StateB>
325  friend auto get_native(const kernel_bundle<StateB> &Obj)
327 
328  template <backend Backend>
330  // NOTE: implementation assumes that the return type is a
331  // derivative of std::vector.
333  ReturnValue.reserve(std::distance(begin(), end()));
334 
335  for (const device_image<State> &DevImg : *this) {
336  ReturnValue.push_back(
337  detail::pi::cast<typename decltype(ReturnValue)::value_type>(
338  DevImg.getNative()));
339  }
340 
341  return ReturnValue;
342  }
343 };
344 #if __cplusplus >= 201703L
345 template <bundle_state State> kernel_bundle(kernel_bundle<State> &&) -> kernel_bundle<State>;
346 #endif
347 
349 // get_kernel_id API
351 
352 namespace detail {
353 // Internal non-template versions of get_kernel_id API which is used by public
354 // onces
355 __SYCL_EXPORT kernel_id get_kernel_id_impl(std::string KernelName);
356 } // namespace detail
357 
359 template <typename KernelName> kernel_id get_kernel_id() {
360  using KI = sycl::detail::KernelInfo<KernelName>;
361  return detail::get_kernel_id_impl(KI::getName());
362 }
363 
365 __SYCL_EXPORT std::vector<kernel_id> get_kernel_ids();
366 
368 // get_kernel_bundle API
370 
371 namespace detail {
372 
373 // Internal non-template versions of get_kernel_bundle API which is used by
374 // public onces
375 __SYCL_EXPORT detail::KernelBundleImplPtr
376 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
377  bundle_state State);
378 
379 __SYCL_EXPORT const std::vector<device>
380 removeDuplicateDevices(const std::vector<device> &Devs);
381 
382 } // namespace detail
383 
388 template <bundle_state State>
390  const std::vector<device> &Devs) {
391  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
392 
394  detail::get_kernel_bundle_impl(Ctx, UniqueDevices, State);
395 
396  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
397 }
398 
399 template <bundle_state State>
401  return get_kernel_bundle<State>(Ctx, Ctx.get_devices());
402 }
403 
404 namespace detail {
405 
406 // Internal non-template versions of get_kernel_bundle API which is used by
407 // public onces
408 __SYCL_EXPORT detail::KernelBundleImplPtr
409 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
410  const std::vector<kernel_id> &KernelIDs,
411  bundle_state State);
412 } // namespace detail
413 
422 template <bundle_state State>
423 kernel_bundle<State>
424 get_kernel_bundle(const context &Ctx, const std::vector<device> &Devs,
425  const std::vector<kernel_id> &KernelIDs) {
426  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
427 
429  detail::get_kernel_bundle_impl(Ctx, UniqueDevices, KernelIDs, State);
430  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
431 }
432 
433 template <bundle_state State>
434 kernel_bundle<State>
435 get_kernel_bundle(const context &Ctx, const std::vector<kernel_id> &KernelIDs) {
436  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(), KernelIDs);
437 }
438 
439 template <typename KernelName, bundle_state State>
441  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(),
442  {get_kernel_id<KernelName>()});
443 }
444 
445 template <typename KernelName, bundle_state State>
447  const std::vector<device> &Devs) {
448  return get_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
449 }
450 
451 namespace detail {
452 
453 // Stable selector function type for passing thru library boundaries
455  std::function<bool(const detail::DeviceImageImplPtr &DevImgImpl)>;
456 
457 // Internal non-template versions of get_kernel_bundle API which is used by
458 // public onces
459 __SYCL_EXPORT detail::KernelBundleImplPtr
460 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
461  bundle_state State, const DevImgSelectorImpl &Selector);
462 
463 // Internal non-template versions of get_empty_interop_kernel_bundle API which
464 // is used by public onces
465 __SYCL_EXPORT detail::KernelBundleImplPtr
467  const std::vector<device> &Devs);
468 
471 template <bundle_state State>
475  return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
476 }
477 } // namespace detail
478 
481 template <bundle_state State, typename SelectorT>
483  const std::vector<device> &Devs,
484  SelectorT Selector) {
485  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
486 
487  detail::DevImgSelectorImpl SelectorWrapper =
488  [Selector](const detail::DeviceImageImplPtr &DevImg) {
489  return Selector(
490  detail::createSyclObjFromImpl<sycl::device_image<State>>(DevImg));
491  };
492 
494  Ctx, UniqueDevices, State, SelectorWrapper);
495 
496  return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
497 }
498 
499 template <bundle_state State, typename SelectorT>
500 kernel_bundle<State> get_kernel_bundle(const context &Ctx, SelectorT Selector) {
501  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(), Selector);
502 }
503 
505 // has_kernel_bundle API
507 
508 namespace detail {
509 
510 __SYCL_EXPORT bool has_kernel_bundle_impl(const context &Ctx,
511  const std::vector<device> &Devs,
512  bundle_state State);
513 
514 __SYCL_EXPORT bool
515 has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
516  const std::vector<kernel_id> &kernelIds,
517  bundle_state State);
518 } // namespace detail
519 
530 template <bundle_state State>
531 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
532  return detail::has_kernel_bundle_impl(Ctx, Devs, State);
533 }
534 
535 template <bundle_state State>
536 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs,
537  const std::vector<kernel_id> &KernelIDs) {
538  return detail::has_kernel_bundle_impl(Ctx, Devs, KernelIDs, State);
539 }
540 
541 template <bundle_state State> bool has_kernel_bundle(const context &Ctx) {
542  return has_kernel_bundle<State>(Ctx, Ctx.get_devices());
543 }
544 
545 template <bundle_state State>
546 bool has_kernel_bundle(const context &Ctx,
547  const std::vector<kernel_id> &KernelIDs) {
548  return has_kernel_bundle<State>(Ctx, Ctx.get_devices(), KernelIDs);
549 }
550 
551 template <typename KernelName, bundle_state State>
552 bool has_kernel_bundle(const context &Ctx) {
553  return has_kernel_bundle<State>(Ctx, {get_kernel_id<KernelName>()});
554 }
555 
556 template <typename KernelName, bundle_state State>
557 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
558  return has_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
559 }
560 
562 // is_compatible API
564 
567 __SYCL_EXPORT bool is_compatible(const std::vector<kernel_id> &KernelIDs,
568  const device &Dev);
569 
570 template <typename KernelName> bool is_compatible(const device &Dev) {
571  return is_compatible({get_kernel_id<KernelName>()}, Dev);
572 }
573 
575 // join API
577 
578 namespace detail {
579 
580 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
581 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
582  bundle_state State);
583 } // namespace detail
584 
587 template <sycl::bundle_state State>
588 sycl::kernel_bundle<State>
589 join(const std::vector<sycl::kernel_bundle<State>> &Bundles) {
590  // Convert kernel_bundle<State> to impls to abstract template parameter away
591  std::vector<detail::KernelBundleImplPtr> KernelBundleImpls;
592  KernelBundleImpls.reserve(Bundles.size());
593  for (const sycl::kernel_bundle<State> &Bundle : Bundles)
594  KernelBundleImpls.push_back(detail::getSyclObjImpl(Bundle));
595 
596  std::shared_ptr<detail::kernel_bundle_impl> Impl =
597  detail::join_impl(KernelBundleImpls, State);
598  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
599 }
600 
602 // compile API
604 
605 namespace detail {
606 
607 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
608 compile_impl(const kernel_bundle<bundle_state::input> &InputBundle,
609  const std::vector<device> &Devs, const property_list &PropList);
610 }
611 
616 inline kernel_bundle<bundle_state::object>
618  const std::vector<device> &Devs, const property_list &PropList = {}) {
619  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
620 
622  detail::compile_impl(InputBundle, UniqueDevices, PropList);
624  kernel_bundle<sycl::bundle_state::object>>(Impl);
625 }
626 
627 inline kernel_bundle<bundle_state::object>
629  const property_list &PropList = {}) {
630  return compile(InputBundle, InputBundle.get_devices(), PropList);
631 }
632 
634 // link API
636 
637 namespace detail {
638 __SYCL_EXPORT std::vector<sycl::device> find_device_intersection(
639  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles);
640 
641 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
642 link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
643  const std::vector<device> &Devs, const property_list &PropList);
644 } // namespace detail
645 
651 inline kernel_bundle<bundle_state::executable>
652 link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
653  const std::vector<device> &Devs, const property_list &PropList = {}) {
654  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
655 
657  detail::link_impl(ObjectBundles, UniqueDevices, PropList);
659  kernel_bundle<sycl::bundle_state::executable>>(Impl);
660 }
661 
662 inline kernel_bundle<bundle_state::executable>
664  const property_list &PropList = {}) {
665  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
666  ObjectBundle.get_devices(), PropList);
667 }
668 
669 inline kernel_bundle<bundle_state::executable>
670 link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
671  const property_list &PropList = {}) {
672  std::vector<sycl::device> IntersectDevices =
673  find_device_intersection(ObjectBundles);
674  return link(ObjectBundles, IntersectDevices, PropList);
675 }
676 
677 inline kernel_bundle<bundle_state::executable>
679  const std::vector<device> &Devs, const property_list &PropList = {}) {
680  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
681  Devs, PropList);
682 }
683 
685 // build API
687 
688 namespace detail {
689 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
690 build_impl(const kernel_bundle<bundle_state::input> &InputBundle,
691  const std::vector<device> &Devs, const property_list &PropList);
692 }
693 
698 inline kernel_bundle<bundle_state::executable>
700  const std::vector<device> &Devs, const property_list &PropList = {}) {
701  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
702 
704  detail::build_impl(InputBundle, UniqueDevices, PropList);
706  kernel_bundle<sycl::bundle_state::executable>>(Impl);
707 }
708 
709 inline kernel_bundle<bundle_state::executable>
711  const property_list &PropList = {}) {
712  return build(InputBundle, InputBundle.get_devices(), PropList);
713 }
714 
715 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
716 } // namespace sycl
717 
718 namespace std {
719 template <> struct hash<sycl::kernel_id> {
720  size_t operator()(const sycl::kernel_id &KernelID) const {
721  return hash<std::shared_ptr<sycl::detail::kernel_id_impl>>()(
722  sycl::detail::getSyclObjImpl(KernelID));
723  }
724 };
725 
726 template <sycl::bundle_state State> struct hash<sycl::device_image<State>> {
727  size_t operator()(const sycl::device_image<State> &DeviceImage) const {
728  return hash<std::shared_ptr<sycl::detail::device_image_impl>>()(
729  sycl::detail::getSyclObjImpl(DeviceImage));
730  }
731 };
732 
733 template <sycl::bundle_state State> struct hash<sycl::kernel_bundle<State>> {
734  size_t operator()(const sycl::kernel_bundle<State> &KernelBundle) const {
735  return hash<std::shared_ptr<sycl::detail::kernel_bundle_impl>>()(
736  sycl::detail::getSyclObjImpl(KernelBundle));
737  }
738 };
739 } // namespace std
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:39
std::vector< device > get_devices() const
Gets devices associated with this SYCL context.
Definition: context.cpp:139
bool operator==(const device_image_plain &RHS) const
device_image_plain(const detail::DeviceImageImplPtr &Impl)
bool operator!=(const device_image_plain &RHS) const
detail::DeviceImageImplPtr impl
pi_native_handle getNative() const
bool operator!=(const kernel_bundle_plain &RHS) const
detail::KernelBundleImplPtr impl
kernel_bundle_plain(const detail::KernelBundleImplPtr &Impl)
bool operator==(const kernel_bundle_plain &RHS) const
Objects of the class represents an instance of an image in a specific state.
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
bool has_kernel(const kernel_id &KernelID) const noexcept
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:47
The kernel_bundle class represents collection of device images in a particular state.
bool contains_specialization_constants() const noexcept
bool empty() const noexcept
device_image_iterator end() const
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
kernel get_kernel(const kernel_id &KernelID) const
bool has_kernel(const kernel_id &KernelID) const noexcept
backend get_backend() const noexcept
bool native_specialization_constant() const noexcept
device_image_iterator begin() const
std::vector< kernel_id > get_kernel_ids() const
friend auto get_native(const kernel_bundle< StateB > &Obj) -> backend_return_t< Backend, kernel_bundle< StateB >>
std::vector< device > get_devices() const noexcept
context get_context() const noexcept
Objects of the class identify kernel is some kernel_bundle related APIs.
bool operator!=(const kernel_id &RHS) const
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:69
Objects of the property_list class are containers for the SYCL properties.
#define __SYCL_INLINE_VER_NAMESPACE(X)
std::shared_ptr< detail::kernel_bundle_impl > join_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
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)
std::shared_ptr< device_image_impl > DeviceImageImplPtr
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
bool has_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, const std::vector< kernel_id > &kernelIds, bundle_state State)
std::vector< sycl::device > find_device_intersection(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles)
detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State, const DevImgSelectorImpl &Selector)
kernel_id get_kernel_id_impl(std::string KernelName)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:240
const std::vector< device > removeDuplicateDevices(const std::vector< device > &Devs)
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: common.hpp:258
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 > KernelBundleImplPtr
kernel_bundle< State > get_empty_interop_kernel_bundle(const context &Ctx)
make_kernel may need an empty interop kernel bundle.
typename std::remove_reference< T >::type remove_reference_t
typename std::enable_if< B, T >::type enable_if_t
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)
sycl::kernel_bundle< State > join(const std::vector< sycl::kernel_bundle< State >> &Bundles)
kernel_bundle< bundle_state::executable > build(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
bool is_compatible(const device &Dev)
kernel_id get_kernel_id()
std::vector< kernel_id > get_kernel_ids()
kernel_bundle< bundle_state::executable > link(const kernel_bundle< bundle_state::object > &ObjectBundle, const std::vector< device > &Devs, const property_list &PropList={})
auto get_native(const SyclObjectT &Obj) -> backend_return_t< BackendName, SyclObjectT >
Definition: backend.hpp:123
float distance(T p0, T p1) __NOEXC
Definition: builtins.hpp:1011
typename backend_traits< Backend >::template return_type< SyclType > backend_return_t
Definition: backend.hpp:72
bool has_kernel_bundle(const context &Ctx, const std::vector< device > &Devs)
kernel_bundle< State > get_kernel_bundle(const context &Ctx, SelectorT Selector)
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const property_list &PropList={})
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
uintptr_t pi_native_handle
Definition: pi.h:107
C++ wrapper of extern "C" PI interfaces.
To cast(From value)
Definition: pi_opencl.cpp:42
size_t operator()(const sycl::device_image< State > &DeviceImage) const
size_t operator()(const sycl::kernel_bundle< State > &KernelBundle) const
size_t operator()(const sycl::kernel_id &KernelID) const