28 template <backend Backend>
class backend_traits;
29 template <backend Backend,
class SyclT>
30 auto get_native(
const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
44 const char *get_name()
const noexcept;
53 kernel_id(
const std::shared_ptr<detail::kernel_id_impl> &Impl)
54 : impl(
std::move(Impl)) {}
56 std::shared_ptr<detail::kernel_id_impl> impl;
66 class device_image_impl;
74 : impl(
std::move(Impl)) {}
77 return impl == RHS.
impl;
81 return !(*
this == RHS);
84 bool has_kernel(
const kernel_id &KernelID)
const noexcept;
86 bool has_kernel(
const kernel_id &KernelID,
const device &Dev)
const noexcept;
102 template <sycl::bundle_state State>
110 return device_image_plain::has_kernel(KernelID);
116 return device_image_plain::has_kernel(KernelID, Dev);
131 class kernel_bundle_impl;
138 : impl(
std::move(Impl)) {}
141 return impl == RHS.
impl;
145 return !(*
this == RHS);
148 bool empty() const noexcept;
150 backend get_backend() const noexcept;
152 context get_context() const noexcept;
154 std::vector<
device> get_devices() const noexcept;
156 bool has_kernel(const
kernel_id &KernelID) const noexcept;
158 bool has_kernel(const
kernel_id &KernelID, const
device &Dev) const noexcept;
162 bool contains_specialization_constants() const noexcept;
164 bool native_specialization_constant() const noexcept;
177 bool has_specialization_constant_impl(const
char *SpecName) const noexcept;
179 void set_specialization_constant_impl(const
char *SpecName,
void *Value,
180 size_t Size) noexcept;
182 void get_specialization_constant_impl(const
char *SpecName,
183 void *Value) const noexcept;
185 bool is_specialization_constant_set(const
char *SpecName) const noexcept;
204 bool empty() const noexcept {
return kernel_bundle_plain::empty(); }
208 return kernel_bundle_plain::get_backend();
213 return kernel_bundle_plain::get_context();
218 return kernel_bundle_plain::get_devices();
224 return kernel_bundle_plain::has_kernel(KernelID);
231 return kernel_bundle_plain::has_kernel(KernelID, Dev);
242 return kernel_bundle_plain::contains_specialization_constants();
248 return kernel_bundle_plain::native_specialization_constant();
256 return detail::kernel_bundle_plain::get_kernel(KernelID);
262 #if __cplusplus >= 201703L
263 template <auto &SpecName>
bool has_specialization_constant() const noexcept {
266 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
267 return has_specialization_constant_impl(SpecSymName);
274 typename = detail::enable_if_t<_State == bundle_state::input>>
275 void set_specialization_constant(
277 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
278 set_specialization_constant_impl(SpecSymName, &Value,
279 sizeof(decltype(Value)));
284 template <auto &SpecName>
286 get_specialization_constant()
const {
287 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
288 if (!is_specialization_constant_set(SpecSymName))
289 return SpecName.getDefaultValue();
294 std::array<
char *,
sizeof(SCType)> RetValue;
296 get_specialization_constant_impl(SpecSymName, RetValue.data());
298 return *
reinterpret_cast<SCType *
>(RetValue.data());
305 kernel_bundle_plain::begin());
313 template <backend Backend>
316 return getNative<Backend>();
329 template <backend Backend,
class SyclT>
330 friend auto get_native(
const SyclT &Obj) -> backend_return_t<Backend, SyclT>;
332 template <backend Backend>
333 backend_return_t<Backend, kernel_bundle<State>> getNative()
const {
336 backend_return_t<Backend, kernel_bundle<State>> ReturnValue;
339 for (
const device_image<State> &DevImg : *
this) {
340 ReturnValue.push_back(
342 DevImg.getNative()));
381 return [](
device a,
device b) {
return a.getNative() != b.getNative(); };
384 inline const std::vector<device>
387 std::set<
device, decltype(compareDevices)> UniqueDeviceSet(
388 Devs.begin(), Devs.end(), compareDevices);
389 std::vector<device> UniqueDevices(UniqueDeviceSet.begin(),
390 UniqueDeviceSet.end());
391 return UniqueDevices;
400 template <bundle_state State>
402 const std::vector<device> &Devs) {
408 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
411 template <bundle_state State>
413 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices());
422 const std::vector<kernel_id> &KernelIDs,
434 template <bundle_state State>
437 const std::vector<kernel_id> &KernelIDs) {
442 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
445 template <bundle_state State>
448 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(), KernelIDs);
451 template <
typename KernelName, bundle_state State>
453 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(),
454 {get_kernel_id<KernelName>()});
457 template <
typename KernelName, bundle_state State>
459 const std::vector<device> &Devs) {
460 return get_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
479 const std::vector<device> &Devs);
483 template <bundle_state State>
487 return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
493 template <bundle_state State,
typename SelectorT>
495 const std::vector<device> &Devs,
496 SelectorT Selector) {
506 Ctx, UniqueDevices, State, SelectorWrapper);
508 return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
511 template <bundle_state State,
typename SelectorT>
513 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(), Selector);
523 const std::vector<device> &Devs,
528 const std::vector<kernel_id> &kernelIds,
542 template <bundle_state State>
547 template <bundle_state State>
549 const std::vector<kernel_id> &KernelIDs) {
554 return has_kernel_bundle<State>(Ctx, Ctx.
get_devices());
557 template <bundle_state State>
559 const std::vector<kernel_id> &KernelIDs) {
560 return has_kernel_bundle<State>(Ctx, Ctx.
get_devices(), KernelIDs);
563 template <
typename KernelName, bundle_state State>
565 return has_kernel_bundle<State>(Ctx, {get_kernel_id<KernelName>()});
568 template <
typename KernelName, bundle_state State>
570 return has_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
579 bool is_compatible(
const std::vector<kernel_id> &KernelIDs,
const device &Dev);
592 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
593 join_impl(
const std::vector<detail::KernelBundleImplPtr> &Bundles);
595 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
596 join_impl(
const std::vector<detail::KernelBundleImplPtr> &Bundles,
602 template <sycl::bundle_state State>
603 sycl::kernel_bundle<State>
606 std::vector<detail::KernelBundleImplPtr> KernelBundleImpls;
607 KernelBundleImpls.reserve(Bundles.size());
611 std::shared_ptr<detail::kernel_bundle_impl> Impl =
613 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
622 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
623 compile_impl(
const kernel_bundle<bundle_state::input> &InputBundle,
624 const std::vector<device> &Devs,
const property_list &PropList);
631 inline kernel_bundle<bundle_state::object>
633 const std::vector<device> &Devs,
const property_list &PropList = {}) {
639 kernel_bundle<sycl::bundle_state::object>>(Impl);
642 inline kernel_bundle<bundle_state::object>
654 const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles);
656 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
657 link_impl(
const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
658 const std::vector<device> &Devs,
const property_list &PropList);
666 inline kernel_bundle<bundle_state::executable>
668 const std::vector<device> &Devs,
const property_list &PropList = {}) {
674 kernel_bundle<sycl::bundle_state::executable>>(Impl);
677 inline kernel_bundle<bundle_state::executable>
680 return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
684 inline kernel_bundle<bundle_state::executable>
687 std::vector<sycl::device> IntersectDevices =
689 return link(ObjectBundles, IntersectDevices, PropList);
692 inline kernel_bundle<bundle_state::executable>
694 const std::vector<device> &Devs,
const property_list &PropList = {}) {
695 return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
704 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
705 build_impl(
const kernel_bundle<bundle_state::input> &InputBundle,
706 const std::vector<device> &Devs,
const property_list &PropList);
713 inline kernel_bundle<bundle_state::executable>
715 const std::vector<device> &Devs,
const property_list &PropList = {}) {
721 kernel_bundle<sycl::bundle_state::executable>>(Impl);
724 inline kernel_bundle<bundle_state::executable>
727 return build(InputBundle, InputBundle.
get_devices(), PropList);
734 template <>
struct hash<
cl::sycl::kernel_id> {
736 return hash<std::shared_ptr<cl::sycl::detail::kernel_id_impl>>()(
741 template <cl::sycl::bundle_state State>
742 struct hash<
cl::sycl::device_image<State>> {
744 return hash<std::shared_ptr<cl::sycl::detail::device_image_impl>>()(
749 template <cl::sycl::bundle_state State>
750 struct hash<
cl::sycl::kernel_bundle<State>> {
752 return hash<std::shared_ptr<cl::sycl::detail::kernel_bundle_impl>>()(