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/backend_types.hpp> // for backend, backend_return_t
12 #include <sycl/context.hpp> // for context
13 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
14 #include <sycl/detail/kernel_desc.hpp> // for get_spec_constant_symboli...
15 #include <sycl/detail/owner_less_base.hpp> // for OwnerLessBase
16 #include <sycl/detail/pi.h> // for pi_native_handle
17 #include <sycl/detail/pi.hpp> // for cast
18 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
20 #endif
21 #include <sycl/device.hpp> // for device
22 #include <sycl/kernel.hpp> // for kernel, kernel_bundle
23 #include <sycl/kernel_bundle_enums.hpp> // for bundle_state
24 #include <sycl/property_list.hpp> // for property_list
25 
26 #include <sycl/ext/oneapi/properties/properties.hpp> // PropertyT
27 #include <sycl/ext/oneapi/properties/property.hpp> // build_options
29 
30 #include <array> // for array
31 #include <cstddef> // for std::byte
32 #include <cstring> // for size_t, memcpy
33 #include <functional> // for function
34 #include <iterator> // for distance
35 #include <memory> // for shared_ptr, operator==, hash
36 #include <string> // for string
37 #include <type_traits> // for enable_if_t, remove_refer...
38 #include <utility> // for move
39 #include <variant> // for hash
40 #include <vector> // for vector
41 
42 namespace sycl {
43 inline namespace _V1 {
44 // Forward declaration
45 template <backend Backend> class backend_traits;
46 template <backend Backend, bundle_state State>
47 auto get_native(const kernel_bundle<State> &Obj)
48  -> backend_return_t<Backend, kernel_bundle<State>>;
49 
50 namespace detail {
51 class kernel_id_impl;
52 class kernel_impl;
53 } // namespace detail
54 
55 template <typename KernelName> kernel_id get_kernel_id();
56 
60 class __SYCL_EXPORT kernel_id : public detail::OwnerLessBase<kernel_id> {
61 public:
62  kernel_id() = delete;
63 
65  const char *get_name() const noexcept;
66 
67  bool operator==(const kernel_id &RHS) const { return impl == RHS.impl; }
68 
69  bool operator!=(const kernel_id &RHS) const { return !(*this == RHS); }
70 
71 private:
72  kernel_id(const char *Name);
73 
74  kernel_id(const std::shared_ptr<detail::kernel_id_impl> &Impl)
75  : impl(std::move(Impl)) {}
76 
77  std::shared_ptr<detail::kernel_id_impl> impl;
78 
79  template <class Obj>
80  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
81 
82  template <class T>
83  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
84 };
85 
86 namespace detail {
87 class device_image_impl;
88 using DeviceImageImplPtr = std::shared_ptr<device_image_impl>;
89 
90 // The class is used as a base for device_image for "untemplating" public
91 // methods.
92 class __SYCL_EXPORT device_image_plain {
93 public:
95  : impl(std::move(Impl)) {}
96 
97  bool operator==(const device_image_plain &RHS) const {
98  return impl == RHS.impl;
99  }
100 
101  bool operator!=(const device_image_plain &RHS) const {
102  return !(*this == RHS);
103  }
104 
105  bool has_kernel(const kernel_id &KernelID) const noexcept;
106 
107  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
108 
109  pi_native_handle getNative() const;
110 
111 protected:
113 
114  template <class Obj>
115  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
116 
117  template <class T>
118  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
119 };
120 } // namespace detail
121 
123 template <sycl::bundle_state State>
125  public detail::OwnerLessBase<device_image<State>> {
126 public:
127  device_image() = delete;
128 
131  bool has_kernel(const kernel_id &KernelID) const noexcept {
132  return device_image_plain::has_kernel(KernelID);
133  }
134 
137  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
138  return device_image_plain::has_kernel(KernelID, Dev);
139  }
140 
141 private:
143  : device_image_plain(std::move(Impl)) {}
144 
145  template <class Obj>
146  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
147 
148  template <class T>
149  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
150 };
151 
152 namespace detail {
153 class kernel_bundle_impl;
154 using KernelBundleImplPtr = std::shared_ptr<detail::kernel_bundle_impl>;
155 
156 // The class is used as a base for kernel_bundle to "untemplate" it's methods
157 class __SYCL_EXPORT kernel_bundle_plain {
158 public:
160  : impl(std::move(Impl)) {}
161 
162  bool operator==(const kernel_bundle_plain &RHS) const {
163  return impl == RHS.impl;
164  }
165 
166  bool operator!=(const kernel_bundle_plain &RHS) const {
167  return !(*this == RHS);
168  }
169 
170  bool empty() const noexcept;
171 
172  backend get_backend() const noexcept;
173 
174  context get_context() const noexcept;
175 
176  std::vector<device> get_devices() const noexcept;
177 
178  bool has_kernel(const kernel_id &KernelID) const noexcept;
179 
180  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept;
181 
182  std::vector<kernel_id> get_kernel_ids() const;
183 
184  bool contains_specialization_constants() const noexcept;
185 
186  bool native_specialization_constant() const noexcept;
187 
188  bool ext_oneapi_has_kernel(const std::string &name);
189 
190  kernel ext_oneapi_get_kernel(const std::string &name);
191 
192 protected:
193  // \returns a kernel object which represents the kernel identified by
194  // kernel_id passed
195  kernel get_kernel(const kernel_id &KernelID) const;
196 
197  // \returns an iterator to the first device image kernel_bundle contains
198  const device_image_plain *begin() const;
199 
200  // \returns an iterator to the last device image kernel_bundle contains
201  const device_image_plain *end() const;
202 
203  bool has_specialization_constant_impl(const char *SpecName) const noexcept;
204 
205  void set_specialization_constant_impl(const char *SpecName, void *Value,
206  size_t Size) noexcept;
207 
208  void get_specialization_constant_impl(const char *SpecName,
209  void *Value) const noexcept;
210 
211  // \returns a bool value which indicates if specialization constant was set to
212  // a value different from default value.
213  bool is_specialization_constant_set(const char *SpecName) const noexcept;
214 
215  detail::KernelBundleImplPtr impl;
216 };
217 
218 } // namespace detail
219 
224 template <bundle_state State>
225 class kernel_bundle : public detail::kernel_bundle_plain,
226  public detail::OwnerLessBase<kernel_bundle<State>> {
227 public:
229 
230  kernel_bundle() = delete;
231 
233  template <
234  bundle_state _State = State,
235  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
236  bool empty() const noexcept {
238  }
239 
243  }
244 
248  }
249 
251  std::vector<device> get_devices() const noexcept {
253  }
254 
257  template <
258  bundle_state _State = State,
259  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
260  bool has_kernel(const kernel_id &KernelID) const noexcept {
261  return kernel_bundle_plain::has_kernel(KernelID);
262  }
263 
267  template <
268  bundle_state _State = State,
269  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
270  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
271  return kernel_bundle_plain::has_kernel(KernelID, Dev);
272  }
273 
276  template <
277  typename KernelName, bundle_state _State = State,
278  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
279  bool has_kernel() const noexcept {
280  return has_kernel(get_kernel_id<KernelName>());
281  }
282 
285  template <
286  typename KernelName, bundle_state _State = State,
287  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
288  bool has_kernel(const device &Dev) const noexcept {
289  return has_kernel(get_kernel_id<KernelName>(), Dev);
290  }
291 
293  template <
294  bundle_state _State = State,
295  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
296  std::vector<kernel_id> get_kernel_ids() const {
298  }
299 
302  template <
303  bundle_state _State = State,
304  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
307  }
308 
311  template <
312  bundle_state _State = State,
313  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
316  }
317 
320  template <bundle_state _State = State,
321  typename = std::enable_if_t<_State == bundle_state::executable>>
322  kernel get_kernel(const kernel_id &KernelID) const {
324  }
325 
328  template <typename KernelName, bundle_state _State = State,
329  typename = std::enable_if_t<_State == bundle_state::executable>>
330  kernel get_kernel() const {
331  return detail::kernel_bundle_plain::get_kernel(get_kernel_id<KernelName>());
332  }
333 
336  template <
337  auto &SpecName, bundle_state _State = State,
338  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
340  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
341  return has_specialization_constant_impl(SpecSymName);
342  }
343 
347  template <auto &SpecName, bundle_state _State = State,
348  typename = std::enable_if_t<_State == bundle_state::input>>
350  typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
351  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
352  set_specialization_constant_impl(SpecSymName, &Value,
353  sizeof(decltype(Value)));
354  }
355 
358  template <
359  auto &SpecName, bundle_state _State = State,
360  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
361  typename std::remove_reference_t<decltype(SpecName)>::value_type
363  using SCType =
364  typename std::remove_reference_t<decltype(SpecName)>::value_type;
365 
366  const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
367  SCType Res{SpecName.getDefaultValue()};
368  if (!is_specialization_constant_set(SpecSymName))
369  return Res;
370 
371  std::array<char, sizeof(SCType)> RetValue;
372  get_specialization_constant_impl(SpecSymName, RetValue.data());
373  std::memcpy(&Res, RetValue.data(), sizeof(SCType));
374 
375  return Res;
376  }
377 
379  template <
380  bundle_state _State = State,
381  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
383  return reinterpret_cast<device_image_iterator>(
385  }
386 
388  template <
389  bundle_state _State = State,
390  typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>>
392  return reinterpret_cast<device_image_iterator>(kernel_bundle_plain::end());
393  }
394 
396  // ext_oneapi_has_kernel
397  // only true if created from source and has this kernel
399  template <bundle_state _State = State,
400  typename = std::enable_if_t<_State == bundle_state::executable>>
401  bool ext_oneapi_has_kernel(const std::string &name) {
403  }
404 
406  // ext_oneapi_get_kernel
407  // kernel_bundle must be created from source, throws if not present
409  template <bundle_state _State = State,
410  typename = std::enable_if_t<_State == bundle_state::executable>>
413  }
414 
415 private:
417  : kernel_bundle_plain(std::move(Impl)) {}
418 
419  template <class Obj>
420  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
421 
422  template <class T>
423  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
424 
425  template <backend Backend, bundle_state StateB>
426  friend auto get_native(const kernel_bundle<StateB> &Obj)
428 
429  template <backend Backend>
431  // NOTE: implementation assumes that the return type is a
432  // derivative of std::vector.
434  ReturnValue.reserve(std::distance(begin(), end()));
435 
436  for (const device_image<State> &DevImg : *this) {
437  ReturnValue.push_back(
438  detail::pi::cast<typename decltype(ReturnValue)::value_type>(
439  DevImg.getNative()));
440  }
441 
442  return ReturnValue;
443  }
444 };
445 template <bundle_state State>
447 
449 // get_kernel_id API
451 
452 namespace detail {
453 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
454 using string_view = std::string;
455 #endif
456 // Internal non-template versions of get_kernel_id API which is used by public
457 // onces
458 __SYCL_EXPORT kernel_id get_kernel_id_impl(string_view KernelName);
459 } // namespace detail
460 
462 template <typename KernelName> kernel_id get_kernel_id() {
463  // FIXME: This must fail at link-time if KernelName not in any available
464  // translation units.
465  using KI = sycl::detail::KernelInfo<KernelName>;
466 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
467  return detail::get_kernel_id_impl(detail::string_view{KI::getName()});
468 #else
469  return detail::get_kernel_id_impl(KI::getName());
470 #endif
471 }
472 
474 __SYCL_EXPORT std::vector<kernel_id> get_kernel_ids();
475 
477 // get_kernel_bundle API
479 
480 namespace detail {
481 
482 // Internal non-template versions of get_kernel_bundle API which is used by
483 // public onces
484 __SYCL_EXPORT detail::KernelBundleImplPtr
485 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
486  bundle_state State);
487 
488 __SYCL_EXPORT const std::vector<device>
489 removeDuplicateDevices(const std::vector<device> &Devs);
490 
491 } // namespace detail
492 
497 template <bundle_state State>
499  const std::vector<device> &Devs) {
500  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
501 
503  detail::get_kernel_bundle_impl(Ctx, UniqueDevices, State);
504 
505  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
506 }
507 
508 template <bundle_state State>
510  return get_kernel_bundle<State>(Ctx, Ctx.get_devices());
511 }
512 
513 namespace detail {
514 
515 // Internal non-template versions of get_kernel_bundle API which is used by
516 // public onces
517 __SYCL_EXPORT detail::KernelBundleImplPtr
518 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
519  const std::vector<kernel_id> &KernelIDs,
520  bundle_state State);
521 } // namespace detail
522 
531 template <bundle_state State>
532 kernel_bundle<State>
533 get_kernel_bundle(const context &Ctx, const std::vector<device> &Devs,
534  const std::vector<kernel_id> &KernelIDs) {
535  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
536 
538  detail::get_kernel_bundle_impl(Ctx, UniqueDevices, KernelIDs, State);
539  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
540 }
541 
542 template <bundle_state State>
543 kernel_bundle<State>
544 get_kernel_bundle(const context &Ctx, const std::vector<kernel_id> &KernelIDs) {
545  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(), KernelIDs);
546 }
547 
548 template <typename KernelName, bundle_state State>
550  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(),
551  {get_kernel_id<KernelName>()});
552 }
553 
554 template <typename KernelName, bundle_state State>
556  const std::vector<device> &Devs) {
557  return get_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
558 }
559 
560 namespace detail {
561 
562 // Stable selector function type for passing thru library boundaries
564  std::function<bool(const detail::DeviceImageImplPtr &DevImgImpl)>;
565 
566 // Internal non-template versions of get_kernel_bundle API which is used by
567 // public onces
568 __SYCL_EXPORT detail::KernelBundleImplPtr
569 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
570  bundle_state State, const DevImgSelectorImpl &Selector);
571 
572 // Internal non-template versions of get_empty_interop_kernel_bundle API which
573 // is used by public onces
574 __SYCL_EXPORT detail::KernelBundleImplPtr
576  const std::vector<device> &Devs);
577 
580 template <bundle_state State>
584  return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
585 }
586 } // namespace detail
587 
590 template <bundle_state State, typename SelectorT>
592  const std::vector<device> &Devs,
593  SelectorT Selector) {
594  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
595 
596  detail::DevImgSelectorImpl SelectorWrapper =
597  [Selector](const detail::DeviceImageImplPtr &DevImg) {
598  return Selector(
600  };
601 
603  Ctx, UniqueDevices, State, SelectorWrapper);
604 
605  return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
606 }
607 
608 template <bundle_state State, typename SelectorT>
609 kernel_bundle<State> get_kernel_bundle(const context &Ctx, SelectorT Selector) {
610  return get_kernel_bundle<State>(Ctx, Ctx.get_devices(), Selector);
611 }
612 
614 // has_kernel_bundle API
616 
617 namespace detail {
618 
619 __SYCL_EXPORT bool has_kernel_bundle_impl(const context &Ctx,
620  const std::vector<device> &Devs,
621  bundle_state State);
622 
623 __SYCL_EXPORT bool
624 has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
625  const std::vector<kernel_id> &kernelIds,
626  bundle_state State);
627 } // namespace detail
628 
639 template <bundle_state State>
640 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
641  return detail::has_kernel_bundle_impl(Ctx, Devs, State);
642 }
643 
644 template <bundle_state State>
645 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs,
646  const std::vector<kernel_id> &KernelIDs) {
647  return detail::has_kernel_bundle_impl(Ctx, Devs, KernelIDs, State);
648 }
649 
650 template <bundle_state State> bool has_kernel_bundle(const context &Ctx) {
651  return has_kernel_bundle<State>(Ctx, Ctx.get_devices());
652 }
653 
654 template <bundle_state State>
655 bool has_kernel_bundle(const context &Ctx,
656  const std::vector<kernel_id> &KernelIDs) {
657  return has_kernel_bundle<State>(Ctx, Ctx.get_devices(), KernelIDs);
658 }
659 
660 template <typename KernelName, bundle_state State>
661 bool has_kernel_bundle(const context &Ctx) {
662  return has_kernel_bundle<State>(Ctx, {get_kernel_id<KernelName>()});
663 }
664 
665 template <typename KernelName, bundle_state State>
666 bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
667  return has_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
668 }
669 
671 // is_compatible API
673 
676 __SYCL_EXPORT bool is_compatible(const std::vector<kernel_id> &KernelIDs,
677  const device &Dev);
678 
679 template <typename KernelName> bool is_compatible(const device &Dev) {
680  return is_compatible({get_kernel_id<KernelName>()}, Dev);
681 }
682 
684 // join API
686 
687 namespace detail {
688 
689 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
690 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
691  bundle_state State);
692 } // namespace detail
693 
696 template <sycl::bundle_state State>
698 join(const std::vector<sycl::kernel_bundle<State>> &Bundles) {
699  // Convert kernel_bundle<State> to impls to abstract template parameter away
700  std::vector<detail::KernelBundleImplPtr> KernelBundleImpls;
701  KernelBundleImpls.reserve(Bundles.size());
702  for (const sycl::kernel_bundle<State> &Bundle : Bundles)
703  KernelBundleImpls.push_back(detail::getSyclObjImpl(Bundle));
704 
705  std::shared_ptr<detail::kernel_bundle_impl> Impl =
706  detail::join_impl(KernelBundleImpls, State);
707  return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
708 }
709 
711 // compile API
713 
714 namespace detail {
715 
716 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
717 compile_impl(const kernel_bundle<bundle_state::input> &InputBundle,
718  const std::vector<device> &Devs, const property_list &PropList);
719 }
720 
725 inline kernel_bundle<bundle_state::object>
727  const std::vector<device> &Devs, const property_list &PropList = {}) {
728  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
729 
731  detail::compile_impl(InputBundle, UniqueDevices, PropList);
733  kernel_bundle<sycl::bundle_state::object>>(Impl);
734 }
735 
736 inline kernel_bundle<bundle_state::object>
738  const property_list &PropList = {}) {
739  return compile(InputBundle, InputBundle.get_devices(), PropList);
740 }
741 
743 // link API
745 
746 namespace detail {
747 __SYCL_EXPORT std::vector<sycl::device> find_device_intersection(
748  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles);
749 
750 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
751 link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
752  const std::vector<device> &Devs, const property_list &PropList);
753 } // namespace detail
754 
760 inline kernel_bundle<bundle_state::executable>
761 link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
762  const std::vector<device> &Devs, const property_list &PropList = {}) {
763  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
764 
766  detail::link_impl(ObjectBundles, UniqueDevices, PropList);
768  kernel_bundle<sycl::bundle_state::executable>>(Impl);
769 }
770 
771 inline kernel_bundle<bundle_state::executable>
773  const property_list &PropList = {}) {
774  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
775  ObjectBundle.get_devices(), PropList);
776 }
777 
778 inline kernel_bundle<bundle_state::executable>
779 link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
780  const property_list &PropList = {}) {
781  std::vector<sycl::device> IntersectDevices =
782  find_device_intersection(ObjectBundles);
783  return link(ObjectBundles, IntersectDevices, PropList);
784 }
785 
786 inline kernel_bundle<bundle_state::executable>
788  const std::vector<device> &Devs, const property_list &PropList = {}) {
789  return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
790  Devs, PropList);
791 }
792 
794 // build API
796 
797 namespace detail {
798 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
799 build_impl(const kernel_bundle<bundle_state::input> &InputBundle,
800  const std::vector<device> &Devs, const property_list &PropList);
801 }
802 
807 inline kernel_bundle<bundle_state::executable>
809  const std::vector<device> &Devs, const property_list &PropList = {}) {
810  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
811 
813  detail::build_impl(InputBundle, UniqueDevices, PropList);
815  kernel_bundle<sycl::bundle_state::executable>>(Impl);
816 }
817 
818 inline kernel_bundle<bundle_state::executable>
820  const property_list &PropList = {}) {
821  return build(InputBundle, InputBundle.get_devices(), PropList);
822 }
823 
824 namespace ext::oneapi::experimental {
825 
827 // PropertyT syclex::build_options
830  : detail::run_time_property_key<detail::PropKind::BuildOptions> {
831  std::vector<std::string> opts;
832  build_options(const std::string &optsArg) : opts{optsArg} {}
833  build_options(const std::vector<std::string> &optsArg) : opts(optsArg) {}
834 };
836 
837 template <>
839  sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
840  : std::true_type {};
841 
843 // PropertyT syclex::save_log
845 struct save_log : detail::run_time_property_key<detail::PropKind::BuildLog> {
847  save_log(std::string *logArg) : log(logArg) {}
848 };
850 
851 template <>
853  sycl::kernel_bundle<bundle_state::ext_oneapi_source>>
854  : std::true_type {};
855 
857 // syclex::is_source_kernel_bundle_supported
859 __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE,
860  source_language Language);
861 
863 // syclex::create_kernel_bundle_from_source
865 
867 create_kernel_bundle_from_source(const context &SyclContext,
868  source_language Language,
869  const std::string &Source);
870 
871 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
873 create_kernel_bundle_from_source(const context &SyclContext,
874  source_language Language,
875  const std::vector<std::byte> &Bytes);
876 #endif
877 
879 // syclex::build(source_kb) => exe_kb
881 namespace detail {
882 // forward decl
885  const std::vector<device> &Devices,
886  const std::vector<std::string> &BuildOptions,
887  std::string *LogPtr);
888 
889 } // namespace detail
890 
891 template <typename PropertyListT = detail::empty_properties_t,
892  typename = std::enable_if_t<
893  is_property_list_v<PropertyListT> &&
896  PropertyListT>::value>>
897 
900  const std::vector<device> &Devices, PropertyListT props = {}) {
901  std::vector<std::string> BuildOptionsVec;
902  std::string *LogPtr = nullptr;
903  if constexpr (props.template has_property<build_options>()) {
904  BuildOptionsVec = props.template get_property<build_options>().opts;
905  }
906  if constexpr (props.template has_property<save_log>()) {
907  LogPtr = props.template get_property<save_log>().log;
908  }
909  return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr);
910 }
911 
912 template <typename PropertyListT = detail::empty_properties_t,
913  typename = std::enable_if_t<
914  is_property_list_v<PropertyListT> &&
915  detail::all_props_are_keys_of<
917  PropertyListT>::value>>
920  PropertyListT props = {}) {
921  return build<PropertyListT>(SourceKB, SourceKB.get_devices(), props);
922 }
923 
924 } // namespace ext::oneapi::experimental
925 
926 } // namespace _V1
927 } // namespace sycl
928 
929 namespace std {
930 template <> struct hash<sycl::kernel_id> {
931  size_t operator()(const sycl::kernel_id &KernelID) const {
932  return hash<std::shared_ptr<sycl::detail::kernel_id_impl>>()(
933  sycl::detail::getSyclObjImpl(KernelID));
934  }
935 };
936 
937 template <sycl::bundle_state State> struct hash<sycl::device_image<State>> {
938  size_t operator()(const sycl::device_image<State> &DeviceImage) const {
939  return hash<std::shared_ptr<sycl::detail::device_image_impl>>()(
940  sycl::detail::getSyclObjImpl(DeviceImage));
941  }
942 };
943 
944 template <sycl::bundle_state State> struct hash<sycl::kernel_bundle<State>> {
945  size_t operator()(const sycl::kernel_bundle<State> &KernelBundle) const {
946  return hash<std::shared_ptr<sycl::detail::kernel_bundle_impl>>()(
947  sycl::detail::getSyclObjImpl(KernelBundle));
948  }
949 };
950 } // namespace std
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:51
std::vector< device > get_devices() const
Gets devices associated with this SYCL context.
Definition: context.cpp:138
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 contains_specialization_constants() const noexcept
std::vector< kernel_id > get_kernel_ids() const
bool operator!=(const kernel_bundle_plain &RHS) const
bool ext_oneapi_has_kernel(const std::string &name)
const device_image_plain * end() const
bool native_specialization_constant() const noexcept
context get_context() const noexcept
detail::KernelBundleImplPtr impl
std::vector< device > get_devices() const noexcept
kernel_bundle_plain(const detail::KernelBundleImplPtr &Impl)
backend get_backend() const noexcept
bool operator==(const kernel_bundle_plain &RHS) const
const device_image_plain * begin() const
bool has_kernel(const kernel_id &KernelID) const noexcept
kernel ext_oneapi_get_kernel(const std::string &name)
kernel get_kernel(const kernel_id &KernelID) 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:65
The kernel_bundle class represents collection of device images in a particular state.
bool ext_oneapi_has_kernel(const std::string &name)
std::vector< kernel_id > get_kernel_ids() const
bool contains_specialization_constants() const noexcept
bool has_kernel(const kernel_id &KernelID) const noexcept
bool has_specialization_constant() const noexcept
device_image_iterator begin() const
device_image_iterator end() const
kernel get_kernel(const kernel_id &KernelID) const
bool has_kernel(const device &Dev) const noexcept
kernel ext_oneapi_get_kernel(const std::string &name)
backend get_backend() const noexcept
friend auto get_native(const kernel_bundle< StateB > &Obj) -> backend_return_t< Backend, kernel_bundle< StateB >>
std::vector< device > get_devices() const noexcept
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
context get_context() const noexcept
bool has_kernel() const noexcept
bool empty() const noexcept
bool native_specialization_constant() const noexcept
std::remove_reference_t< decltype(SpecName)>::value_type get_specialization_constant() const
void set_specialization_constant(typename std::remove_reference_t< decltype(SpecName)>::value_type Value)
Sets the value of the specialization constant whose address is SpecName for this bundle.
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:74
Objects of the property_list class are containers for the SYCL properties.
PiProgram cast(cl_program)=delete
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)
std::shared_ptr< device_image_impl > DeviceImageImplPtr
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
std::string string_view
Definition: handler.hpp:424
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)
T createSyclObjFromImpl(decltype(T::impl) ImplObj)
Definition: impl_utils.hpp:48
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
kernel_bundle< State > get_empty_interop_kernel_bundle(const context &Ctx)
make_kernel may need an empty interop kernel bundle.
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)
kernel_bundle< bundle_state::executable > build_from_source(kernel_bundle< bundle_state::ext_oneapi_source > &SourceKB, const std::vector< device > &Devices, const std::vector< std::string > &BuildOptions, std::string *LogPtr)
kernel_bundle< bundle_state::ext_oneapi_source > create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::vector< std::byte > &Bytes)
bool is_source_kernel_bundle_supported(backend BE, source_language Language)
sycl::detail::kernel_bundle_impl kernel_bundle_impl
decltype(properties{}) empty_properties_t
Definition: properties.hpp:190
kernel_bundle(kernel_bundle< State > &&) -> kernel_bundle< State >
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={})
float distance(T p0, T p1)
kernel_id get_kernel_id()
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
std::vector< kernel_id > get_kernel_ids()
bool is_compatible(const std::vector< kernel_id > &KernelIDs, const device &Dev)
bool has_kernel_bundle(const context &Ctx, const std::vector< device > &Devs)
kernel_bundle< State > get_kernel_bundle(const context &Ctx, const std::vector< device > &Devs)
A kernel bundle in state State which contains all of the kernels in the application which are compati...
ESIMD_NODEBUG ESIMD_INLINE sycl::ext::intel::esimd::simd< float, SZ > log(sycl::ext::intel::esimd::simd< float, SZ > x) __NOEXC
auto get_native(const SyclObjectT &Obj) -> backend_return_t< BackendName, SyclObjectT >
Definition: backend.hpp:136
typename backend_traits< Backend >::template return_type< SyclType > backend_return_t
Definition: backend.hpp:87
kernel_bundle< bundle_state::executable > link(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, const std::vector< device > &Devs, const property_list &PropList={})
Definition: access.hpp:18
uintptr_t pi_native_handle
Definition: pi.h:206
C++ wrapper of extern "C" PI interfaces.
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
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
build_options(const std::vector< std::string > &optsArg)