30 template <backend Backend>
class backend_traits;
31 template <backend Backend, bundle_state State>
32 auto get_native(
const kernel_bundle<State> &Obj)
33 -> backend_return_t<Backend, kernel_bundle<State>>;
49 const char *get_name()
const noexcept;
58 kernel_id(
const std::shared_ptr<detail::kernel_id_impl> &Impl)
59 : impl(
std::move(Impl)) {}
61 std::shared_ptr<detail::kernel_id_impl> impl;
71 class device_image_impl;
79 : impl(
std::move(Impl)) {}
82 return impl == RHS.
impl;
86 return !(*
this == RHS);
89 bool has_kernel(
const kernel_id &KernelID)
const noexcept;
91 bool has_kernel(
const kernel_id &KernelID,
const device &Dev)
const noexcept;
107 template <sycl::bundle_state State>
116 return device_image_plain::has_kernel(KernelID);
122 return device_image_plain::has_kernel(KernelID, Dev);
127 : device_image_plain(
std::move(Impl)) {}
137 class kernel_bundle_impl;
144 : impl(
std::move(Impl)) {}
147 return impl == RHS.
impl;
151 return !(*
this == RHS);
154 bool empty() const noexcept;
156 backend get_backend() const noexcept;
158 context get_context() const noexcept;
160 std::vector<
device> get_devices() const noexcept;
162 bool has_kernel(const
kernel_id &KernelID) const noexcept;
164 bool has_kernel(const
kernel_id &KernelID, const
device &Dev) const noexcept;
168 bool contains_specialization_constants() const noexcept;
170 bool native_specialization_constant() const noexcept;
183 bool has_specialization_constant_impl(const
char *SpecName) const noexcept;
185 void set_specialization_constant_impl(const
char *SpecName,
void *Value,
186 size_t Size) noexcept;
188 void get_specialization_constant_impl(const
char *SpecName,
189 void *Value) const noexcept;
191 bool is_specialization_constant_set(const
char *SpecName) const noexcept;
211 bool empty() const noexcept {
return kernel_bundle_plain::empty(); }
215 return kernel_bundle_plain::get_backend();
220 return kernel_bundle_plain::get_context();
225 return kernel_bundle_plain::get_devices();
231 return kernel_bundle_plain::has_kernel(KernelID);
238 return kernel_bundle_plain::has_kernel(KernelID, Dev);
243 template <
typename KernelName>
bool has_kernel() const noexcept {
244 return has_kernel(get_kernel_id<KernelName>());
249 template <
typename KernelName>
251 return has_kernel(get_kernel_id<KernelName>(), Dev);
262 return kernel_bundle_plain::contains_specialization_constants();
268 return kernel_bundle_plain::native_specialization_constant();
276 return detail::kernel_bundle_plain::get_kernel(KernelID);
281 template <
typename KernelName,
bundle_state _State = State,
284 return detail::kernel_bundle_plain::get_kernel(get_kernel_id<KernelName>());
290 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
291 return has_specialization_constant_impl(SpecSymName);
301 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
302 set_specialization_constant_impl(SpecSymName, &Value,
303 sizeof(decltype(Value)));
308 template <auto &SpecName>
314 const char *SpecSymName = detail::get_spec_constant_symbolic_ID<SpecName>();
315 SCType Res{SpecName.getDefaultValue()};
316 if (!is_specialization_constant_set(SpecSymName))
319 std::array<char,
sizeof(SCType)> RetValue;
320 get_specialization_constant_impl(SpecSymName, RetValue.data());
321 std::memcpy(&Res, RetValue.data(),
sizeof(SCType));
329 kernel_bundle_plain::begin());
339 : kernel_bundle_plain(
std::move(Impl)) {}
347 template <backend Backend, bundle_state StateB>
348 friend auto get_native(
const kernel_bundle<StateB> &Obj)
349 -> backend_return_t<Backend, kernel_bundle<StateB>>;
351 template <backend Backend>
352 backend_return_t<Backend, kernel_bundle<State>> getNative()
const {
355 backend_return_t<Backend, kernel_bundle<State>> ReturnValue;
356 ReturnValue.reserve(std::distance(begin(), end()));
358 for (
const device_image<State> &DevImg : *
this) {
359 ReturnValue.push_back(
361 DevImg.getNative()));
367 template <bundle_state State>
368 kernel_bundle(kernel_bundle<State> &&) -> kernel_bundle<State>;
384 using KI = sycl::detail::KernelInfo<KernelName>;
403 __SYCL_EXPORT
const std::vector<device>
412 template <bundle_state State>
414 const std::vector<device> &Devs) {
420 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
423 template <bundle_state State>
425 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices());
434 const std::vector<kernel_id> &KernelIDs,
446 template <bundle_state State>
449 const std::vector<kernel_id> &KernelIDs) {
454 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
457 template <bundle_state State>
460 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(), KernelIDs);
463 template <
typename KernelName, bundle_state State>
465 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(),
466 {get_kernel_id<KernelName>()});
469 template <
typename KernelName, bundle_state State>
471 const std::vector<device> &Devs) {
472 return get_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
491 const std::vector<device> &Devs);
495 template <bundle_state State>
499 return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
505 template <bundle_state State,
typename SelectorT>
507 const std::vector<device> &Devs,
508 SelectorT Selector) {
518 Ctx, UniqueDevices, State, SelectorWrapper);
520 return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
523 template <bundle_state State,
typename SelectorT>
525 return get_kernel_bundle<State>(Ctx, Ctx.
get_devices(), Selector);
535 const std::vector<device> &Devs,
540 const std::vector<kernel_id> &kernelIds,
554 template <bundle_state State>
559 template <bundle_state State>
561 const std::vector<kernel_id> &KernelIDs) {
566 return has_kernel_bundle<State>(Ctx, Ctx.
get_devices());
569 template <bundle_state State>
571 const std::vector<kernel_id> &KernelIDs) {
572 return has_kernel_bundle<State>(Ctx, Ctx.
get_devices(), KernelIDs);
575 template <
typename KernelName, bundle_state State>
577 return has_kernel_bundle<State>(Ctx, {get_kernel_id<KernelName>()});
580 template <
typename KernelName, bundle_state State>
582 return has_kernel_bundle<State>(Ctx, Devs, {get_kernel_id<KernelName>()});
591 __SYCL_EXPORT
bool is_compatible(
const std::vector<kernel_id> &KernelIDs,
604 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
605 join_impl(
const std::vector<detail::KernelBundleImplPtr> &Bundles,
611 template <sycl::bundle_state State>
612 sycl::kernel_bundle<State>
613 join(
const std::vector<sycl::kernel_bundle<State>> &Bundles) {
615 std::vector<detail::KernelBundleImplPtr> KernelBundleImpls;
616 KernelBundleImpls.reserve(Bundles.size());
617 for (
const sycl::kernel_bundle<State> &Bundle : Bundles)
620 std::shared_ptr<detail::kernel_bundle_impl> Impl =
622 return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
631 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
632 compile_impl(
const kernel_bundle<bundle_state::input> &InputBundle,
633 const std::vector<device> &Devs,
const property_list &PropList);
640 inline kernel_bundle<bundle_state::object>
642 const std::vector<device> &Devs,
const property_list &PropList = {}) {
648 kernel_bundle<sycl::bundle_state::object>>(Impl);
651 inline kernel_bundle<bundle_state::object>
663 const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles);
665 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
666 link_impl(
const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
667 const std::vector<device> &Devs,
const property_list &PropList);
675 inline kernel_bundle<bundle_state::executable>
677 const std::vector<device> &Devs,
const property_list &PropList = {}) {
683 kernel_bundle<sycl::bundle_state::executable>>(Impl);
686 inline kernel_bundle<bundle_state::executable>
689 return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
693 inline kernel_bundle<bundle_state::executable>
696 std::vector<sycl::device> IntersectDevices =
698 return link(ObjectBundles, IntersectDevices, PropList);
701 inline kernel_bundle<bundle_state::executable>
703 const std::vector<device> &Devs,
const property_list &PropList = {}) {
704 return link(std::vector<kernel_bundle<bundle_state::object>>{ObjectBundle},
713 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
714 build_impl(
const kernel_bundle<bundle_state::input> &InputBundle,
715 const std::vector<device> &Devs,
const property_list &PropList);
722 inline kernel_bundle<bundle_state::executable>
724 const std::vector<device> &Devs,
const property_list &PropList = {}) {
730 kernel_bundle<sycl::bundle_state::executable>>(Impl);
733 inline kernel_bundle<bundle_state::executable>
736 return build(InputBundle, InputBundle.
get_devices(), PropList);
743 template <>
struct hash<
sycl::kernel_id> {
745 return hash<std::shared_ptr<sycl::detail::kernel_id_impl>>()(
750 template <sycl::bundle_state State>
struct hash<
sycl::device_image<State>> {
751 size_t operator()(
const sycl::device_image<State> &DeviceImage)
const {
752 return hash<std::shared_ptr<sycl::detail::device_image_impl>>()(
758 size_t operator()(
const sycl::kernel_bundle<State> &KernelBundle)
const {
759 return hash<std::shared_ptr<sycl::detail::kernel_bundle_impl>>()(