18 inline namespace _V1 {
21 : impl(
std::make_shared<detail::kernel_id_impl>(Name)) {}
23 const char *kernel_id::get_name() const
noexcept {
return impl->get_name(); }
32 return impl->has_kernel(KernelID);
35 bool device_image_plain::has_kernel(
const kernel_id &KernelID,
37 return impl->has_kernel(KernelID, Dev);
41 return impl->getNative();
48 bool kernel_bundle_plain::empty() const
noexcept {
return impl->empty(); }
51 return impl->get_backend();
55 return impl->get_context();
58 std::vector<device> kernel_bundle_plain::get_devices() const
noexcept {
59 return impl->get_devices();
63 return impl->get_kernel_ids();
66 bool kernel_bundle_plain::contains_specialization_constants() const
noexcept {
67 return impl->contains_specialization_constants();
70 bool kernel_bundle_plain::native_specialization_constant() const
noexcept {
71 return impl->native_specialization_constant();
75 return impl->get_kernel(KernelID, impl);
87 return impl->has_kernel(KernelID);
90 bool kernel_bundle_plain::has_kernel(
const kernel_id &KernelID,
92 return impl->has_kernel(KernelID, Dev);
95 bool kernel_bundle_plain::has_specialization_constant_impl(
96 const char *SpecName)
const noexcept {
97 return impl->has_specialization_constant(SpecName);
100 void kernel_bundle_plain::set_specialization_constant_impl(
101 const char *SpecName,
void *Value,
size_t Size)
noexcept {
102 impl->set_specialization_constant_raw_value(SpecName, Value, Size);
105 void kernel_bundle_plain::get_specialization_constant_impl(
106 const char *SpecName,
void *Value)
const noexcept {
107 impl->get_specialization_constant_raw_value(SpecName, Value);
110 bool kernel_bundle_plain::is_specialization_constant_set(
111 const char *SpecName)
const noexcept {
112 return impl->is_specialization_constant_set(SpecName);
115 bool kernel_bundle_plain::ext_oneapi_has_kernel(
const std::string &name) {
116 return impl->ext_oneapi_has_kernel(name);
119 kernel kernel_bundle_plain::ext_oneapi_get_kernel(
const std::string &name) {
120 return impl->ext_oneapi_get_kernel(name, impl);
127 const std::vector<device>
129 std::vector<device> UniqueDevices;
132 std::unordered_set<device> UniqueDeviceSet;
133 for (
const device &Dev : Devs)
134 if (UniqueDeviceSet.insert(Dev).second)
135 UniqueDevices.push_back(Dev);
137 return UniqueDevices;
141 return detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
147 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, State);
152 const std::vector<kernel_id> &KernelIDs,
154 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, KernelIDs,
161 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, Selector,
167 const std::vector<device> &Devs) {
168 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
171 std::shared_ptr<detail::kernel_bundle_impl>
172 join_impl(
const std::vector<detail::KernelBundleImplPtr> &Bundles,
174 return std::make_shared<detail::kernel_bundle_impl>(Bundles, State);
181 if (Devs.empty() || !AllDevicesInTheContext)
183 "Not all devices are associated with the context or "
184 "vector of devices is empty");
186 if (bundle_state::input == State &&
189 if (bundle_state::object == State &&
193 const std::vector<device_image_plain> DeviceImages =
194 detail::ProgramManager::getInstance()
195 .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
197 return (
bool)DeviceImages.size();
201 const std::vector<kernel_id> &KernelIds,
206 if (Devs.empty() || !AllDevicesInTheContext)
208 "Not all devices are associated with the context or "
209 "vector of devices is empty");
211 bool DeviceHasRequireAspectForState =
true;
212 if (bundle_state::input == State) {
213 DeviceHasRequireAspectForState =
215 return Dev.has(aspect::online_compiler);
217 }
else if (bundle_state::object == State) {
218 DeviceHasRequireAspectForState =
220 return Dev.has(aspect::online_linker);
224 if (!DeviceHasRequireAspectForState)
227 const std::vector<device_image_plain> DeviceImages =
228 detail::ProgramManager::getInstance()
229 .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
231 std::set<kernel_id, LessByNameComp> CombinedKernelIDs;
233 const std::shared_ptr<device_image_impl> &DeviceImageImpl =
236 CombinedKernelIDs.insert(DeviceImageImpl->get_kernel_ids_ptr()->begin(),
237 DeviceImageImpl->get_kernel_ids_ptr()->end());
240 const bool AllKernelIDsRepresented =
242 [&CombinedKernelIDs](
const kernel_id &KernelID) {
243 return CombinedKernelIDs.count(KernelID);
246 return AllKernelIDsRepresented;
249 std::shared_ptr<detail::kernel_bundle_impl>
251 const std::vector<device> &Devs,
const property_list &PropList) {
252 return std::make_shared<detail::kernel_bundle_impl>(
253 InputBundle, Devs, PropList, bundle_state::object);
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>(ObjectBundles, Devs,
263 std::shared_ptr<detail::kernel_bundle_impl>
265 const std::vector<device> &Devs,
const property_list &PropList) {
266 return std::make_shared<detail::kernel_bundle_impl>(
267 InputBundle, Devs, PropList, bundle_state::executable);
274 std::vector<sycl::device> IntersectDevices;
275 std::vector<unsigned int> DevsCounters;
276 std::map<device, unsigned int, LessByHash<device>> DevCounters;
281 DevCounters[Device]++;
285 for (
const std::pair<const device, unsigned int> &It : DevCounters)
286 if (ObjectBundles.size() == It.second)
287 IntersectDevices.push_back(It.first);
289 return IntersectDevices;
299 return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
303 if (KernelIDs.empty())
308 auto doesImageTargetMatchDevice = [](
const device &Dev,
310 const char *Target = Img.getRawData().DeviceTargetSpec;
313 if (BE == sycl::backend::ext_intel_esimd_emulator) {
318 return (BE == sycl::backend::opencl ||
319 BE == sycl::backend::ext_oneapi_level_zero);
325 return Dev.
is_gpu() && (BE == sycl::backend::opencl ||
326 BE == sycl::backend::ext_oneapi_level_zero);
331 return BE == sycl::backend::ext_oneapi_cuda;
333 return BE == sycl::backend::ext_oneapi_hip;
343 for (
const auto &KernelID : KernelIDs) {
344 std::set<detail::RTDeviceBinaryImage *> BinImages =
345 detail::ProgramManager::getInstance().getRawDeviceImages({KernelID});
349 return doesDevSupportDeviceRequirements(Dev, *Img) &&
350 doesImageTargetMatchDevice(Dev, *Img);
361 namespace ext::oneapi::experimental {
373 bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
374 (BE == sycl::backend::opencl);
375 if ((Language == source_language::opencl) && BE_Acceptable) {
388 const std::string &Source) {
395 "kernel_bundle creation from source not supported");
397 std::shared_ptr<kernel_bundle_impl> KBImpl =
398 std::make_shared<kernel_bundle_impl>(SyclContext, Language, Source);
399 return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
408 const std::vector<device> &Devices,
410 std::string *LogPtr) {
411 std::vector<device> UniqueDevices =
413 std::shared_ptr<kernel_bundle_impl> sourceImpl =
getSyclObjImpl(SourceKB);
414 std::shared_ptr<kernel_bundle_impl> KBImpl =
415 sourceImpl->build_from_source(UniqueDevices,
BuildOptions, LogPtr);
416 return sycl::detail::createSyclObjFromImpl<exe_kb>(KBImpl);