DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_bundle_impl.hpp
Go to the documentation of this file.
1 //==------- kernel_bundle_impl.hpp - SYCL kernel_bundle_impl ---------------==//
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 
13 #include <detail/kernel_impl.hpp>
15 #include <sycl/backend_types.hpp>
16 #include <sycl/context.hpp>
17 #include <sycl/detail/common.hpp>
19 #include <sycl/detail/pi.h>
20 #include <sycl/device.hpp>
21 #include <sycl/kernel_bundle.hpp>
22 
23 #include <algorithm>
24 #include <cassert>
25 #include <cstdint>
26 #include <cstring>
27 #include <memory>
28 #include <vector>
29 
30 namespace sycl {
31 inline namespace _V1 {
32 namespace detail {
33 
34 static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,
35  const context &Context) {
36  return std::all_of(
37  Devices.begin(), Devices.end(), [&Context](const device &Dev) {
38  return getSyclObjImpl(Context)->isDeviceValid(getSyclObjImpl(Dev));
39  });
40 }
41 
42 static bool checkAllDevicesHaveAspect(const std::vector<device> &Devices,
43  aspect Aspect) {
44  return std::all_of(Devices.begin(), Devices.end(),
45  [&Aspect](const device &Dev) { return Dev.has(Aspect); });
46 }
47 
48 namespace syclex = sycl::ext::oneapi::experimental;
49 
50 class kernel_impl;
51 
53 // It provides an access and utilities to manage set of sycl::device_images
54 // objects.
56 
57  using SpecConstMapT = std::map<std::string, std::vector<unsigned char>>;
58 
59  void common_ctor_checks(bundle_state State) {
60  const bool AllDevicesInTheContext =
61  checkAllDevicesAreInContext(MDevices, MContext);
62  if (MDevices.empty() || !AllDevicesInTheContext)
63  throw sycl::exception(
65  "Not all devices are associated with the context or "
66  "vector of devices is empty");
67 
68  if (bundle_state::input == State &&
69  !checkAllDevicesHaveAspect(MDevices, aspect::online_compiler))
71  "Not all devices have aspect::online_compiler");
72 
73  if (bundle_state::object == State &&
74  !checkAllDevicesHaveAspect(MDevices, aspect::online_linker))
76  "Not all devices have aspect::online_linker");
77  }
78 
79 public:
80  kernel_bundle_impl(context Ctx, std::vector<device> Devs, bundle_state State)
81  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
82 
83  common_ctor_checks(State);
84 
86  MContext, MDevices, State);
87  }
88 
89  // Interop constructor used by make_kernel
90  kernel_bundle_impl(context Ctx, std::vector<device> Devs)
91  : MContext(Ctx), MDevices(Devs), MState(bundle_state::executable) {
92  if (!checkAllDevicesAreInContext(Devs, Ctx))
93  throw sycl::exception(
95  "Not all devices are associated with the context or "
96  "vector of devices is empty");
97  MIsInterop = true;
98  }
99 
100  // Interop constructor
101  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
102  device_image_plain &DevImage)
103  : kernel_bundle_impl(Ctx, Devs) {
104  MDeviceImages.push_back(DevImage);
105  }
106 
107  // Matches sycl::build and sycl::compile
108  // Have one constructor because sycl::build and sycl::compile have the same
109  // signature
111  std::vector<device> Devs, const property_list &PropList,
112  bundle_state TargetState)
113  : MContext(InputBundle.get_context()), MDevices(std::move(Devs)),
114  MState(TargetState) {
115 
116  MSpecConstValues = getSyclObjImpl(InputBundle)->get_spec_const_map_ref();
117 
118  const std::vector<device> &InputBundleDevices =
119  getSyclObjImpl(InputBundle)->get_devices();
120  const bool AllDevsAssociatedWithInputBundle =
121  std::all_of(MDevices.begin(), MDevices.end(),
122  [&InputBundleDevices](const device &Dev) {
123  return InputBundleDevices.end() !=
124  std::find(InputBundleDevices.begin(),
125  InputBundleDevices.end(), Dev);
126  });
127  if (MDevices.empty() || !AllDevsAssociatedWithInputBundle)
128  throw sycl::exception(
130  "Not all devices are in the set of associated "
131  "devices for input bundle or vector of devices is empty");
132 
133  for (const device_image_plain &DeviceImage : InputBundle) {
134  // Skip images which are not compatible with devices provided
135  if (std::none_of(
136  MDevices.begin(), MDevices.end(),
137  [&DeviceImage](const device &Dev) {
138  return getSyclObjImpl(DeviceImage)->compatible_with_device(Dev);
139  }))
140  continue;
141 
142  switch (TargetState) {
144  MDeviceImages.push_back(detail::ProgramManager::getInstance().compile(
145  DeviceImage, MDevices, PropList));
146  break;
148  MDeviceImages.push_back(detail::ProgramManager::getInstance().build(
149  DeviceImage, MDevices, PropList));
150  break;
151  case bundle_state::input:
153  throw sycl::runtime_error("Internal error. The target state should not "
154  "be input or ext_oneapi_source",
155  PI_ERROR_INVALID_OPERATION);
156  break;
157  }
158  }
159  }
160 
161  // Matches sycl::link
163  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
164  std::vector<device> Devs, const property_list &PropList)
165  : MDevices(std::move(Devs)), MState(bundle_state::executable) {
166 
167  if (MDevices.empty())
169  "Vector of devices is empty");
170 
171  if (ObjectBundles.empty())
172  return;
173 
174  MContext = ObjectBundles[0].get_context();
175  for (size_t I = 1; I < ObjectBundles.size(); ++I) {
176  if (ObjectBundles[I].get_context() != MContext)
177  throw sycl::exception(
179  "Not all input bundles have the same associated context");
180  }
181 
182  // Check if any of the devices in devs are not in the set of associated
183  // devices for any of the bundles in ObjectBundles
184  const bool AllDevsAssociatedWithInputBundles = std::all_of(
185  MDevices.begin(), MDevices.end(), [&ObjectBundles](const device &Dev) {
186  // Number of devices is expected to be small
187  return std::all_of(
188  ObjectBundles.begin(), ObjectBundles.end(),
189  [&Dev](const kernel_bundle<bundle_state::object> &KernelBundle) {
190  const std::vector<device> &BundleDevices =
191  getSyclObjImpl(KernelBundle)->get_devices();
192  return BundleDevices.end() != std::find(BundleDevices.begin(),
193  BundleDevices.end(),
194  Dev);
195  });
196  });
197  if (!AllDevsAssociatedWithInputBundles)
199  "Not all devices are in the set of associated "
200  "devices for input bundles");
201 
202  // TODO: Unify with c'tor for sycl::comile and sycl::build by calling
203  // sycl::join on vector of kernel_bundles
204 
205  // The loop below just links each device image separately, not linking any
206  // two device images together. This is correct so long as each device image
207  // has no unresolved symbols. That's the case when device images are created
208  // from generic SYCL APIs. There's no way in generic SYCL to create a kernel
209  // which references an undefined symbol. If we decide in the future to allow
210  // a backend interop API to create a "sycl::kernel_bundle" that references
211  // undefined symbols, then the logic in this loop will need to be changed.
212  for (const kernel_bundle<bundle_state::object> &ObjectBundle :
213  ObjectBundles) {
214  for (const device_image_plain &DeviceImage : ObjectBundle) {
215 
216  // Skip images which are not compatible with devices provided
217  if (std::none_of(MDevices.begin(), MDevices.end(),
218  [&DeviceImage](const device &Dev) {
219  return getSyclObjImpl(DeviceImage)
220  ->compatible_with_device(Dev);
221  }))
222  continue;
223 
224  std::vector<device_image_plain> LinkedResults =
225  detail::ProgramManager::getInstance().link(DeviceImage, MDevices,
226  PropList);
227  MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
228  LinkedResults.end());
229  }
230  }
231 
232  for (const kernel_bundle<bundle_state::object> &Bundle : ObjectBundles) {
233  const KernelBundleImplPtr BundlePtr = getSyclObjImpl(Bundle);
234  for (const std::pair<const std::string, std::vector<unsigned char>>
235  &SpecConst : BundlePtr->MSpecConstValues) {
236  MSpecConstValues[SpecConst.first] = SpecConst.second;
237  }
238  }
239  }
240 
241  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
242  const std::vector<kernel_id> &KernelIDs,
243  bundle_state State)
244  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
245 
246  common_ctor_checks(State);
247 
249  MContext, MDevices, KernelIDs, State);
250  }
251 
252  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
253  const DevImgSelectorImpl &Selector, bundle_state State)
254  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
255 
256  common_ctor_checks(State);
257 
259  MContext, MDevices, Selector, State);
260  }
261 
262  // C'tor matches sycl::join API
263  kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
264  bundle_state State)
265  : MState(State) {
266  if (Bundles.empty())
267  return;
268 
269  MContext = Bundles[0]->MContext;
270  MDevices = Bundles[0]->MDevices;
271  for (size_t I = 1; I < Bundles.size(); ++I) {
272  if (Bundles[I]->MContext != MContext)
273  throw sycl::exception(
275  "Not all input bundles have the same associated context.");
276  if (Bundles[I]->MDevices != MDevices)
277  throw sycl::exception(
279  "Not all input bundles have the same set of associated devices.");
280  }
281 
282  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
283 
284  MDeviceImages.insert(MDeviceImages.end(), Bundle->MDeviceImages.begin(),
285  Bundle->MDeviceImages.end());
286  }
287 
288  std::sort(MDeviceImages.begin(), MDeviceImages.end(),
290 
291  if (get_bundle_state() == bundle_state::input) {
292  // Copy spec constants values from the device images to be removed.
293  auto MergeSpecConstants = [this](const device_image_plain &Img) {
294  const detail::DeviceImageImplPtr &ImgImpl = getSyclObjImpl(Img);
295  const std::map<std::string,
296  std::vector<device_image_impl::SpecConstDescT>>
297  &SpecConsts = ImgImpl->get_spec_const_data_ref();
298  const std::vector<unsigned char> &Blob =
299  ImgImpl->get_spec_const_blob_ref();
300  for (const std::pair<const std::string,
301  std::vector<device_image_impl::SpecConstDescT>>
302  &SpecConst : SpecConsts) {
303  if (SpecConst.second.front().IsSet)
304  set_specialization_constant_raw_value(
305  SpecConst.first.c_str(),
306  Blob.data() + SpecConst.second.front().BlobOffset,
307  SpecConst.second.back().CompositeOffset +
308  SpecConst.second.back().Size);
309  }
310  };
311  std::for_each(MDeviceImages.begin(), MDeviceImages.end(),
312  MergeSpecConstants);
313  }
314 
315  const auto DevImgIt =
316  std::unique(MDeviceImages.begin(), MDeviceImages.end());
317 
318  // Remove duplicate device images.
319  MDeviceImages.erase(DevImgIt, MDeviceImages.end());
320 
321  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
322  for (const std::pair<const std::string, std::vector<unsigned char>>
323  &SpecConst : Bundle->MSpecConstValues) {
324  set_specialization_constant_raw_value(SpecConst.first.c_str(),
325  SpecConst.second.data(),
326  SpecConst.second.size());
327  }
328  }
329  }
330 
331  // oneapi_ext_kernel_compiler
332  // construct from source string
334  const std::string &Src)
335  : MContext(Context), MDevices(Context.get_devices()),
336  MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {}
337 
338  // oneapi_ext_kernel_compiler
339  // construct from source bytes
341  const std::vector<std::byte> &Bytes)
342  : MContext(Context), MDevices(Context.get_devices()),
343  MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Bytes) {
344  }
345 
346  // oneapi_ext_kernel_compiler
347  // interop constructor
348  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
349  device_image_plain &DevImage,
350  std::vector<std::string> KNames)
351  : kernel_bundle_impl(Ctx, Devs, DevImage) {
352  MState = bundle_state::executable;
353  KernelNames = KNames;
354  }
355 
356  std::shared_ptr<kernel_bundle_impl>
357  build_from_source(const std::vector<device> Devices,
358  const std::vector<std::string> &BuildOptions,
359  std::string *LogPtr) {
360  assert(MState == bundle_state::ext_oneapi_source &&
361  "bundle_state::ext_oneapi_source required");
362 
363  const auto spirv = [&]() -> std::vector<uint8_t> {
364  if (Language == syclex::source_language::opencl) {
365  // if successful, the log is empty. if failed, throws an error with the
366  // compilation log.
367  const auto &SourceStr = std::get<std::string>(this->Source);
369  LogPtr);
370  }
371  if (Language == syclex::source_language::spirv) {
372  const auto &SourceBytes =
373  std::get<std::vector<std::byte>>(this->Source);
374  std::vector<uint8_t> Result(SourceBytes.size());
375  std::transform(SourceBytes.cbegin(), SourceBytes.cend(), Result.begin(),
376  [](std::byte B) { return static_cast<uint8_t>(B); });
377  return Result;
378  }
379  throw sycl::exception(
381  "OpenCL C and SPIR-V are the only supported languages at this time");
382  }();
383 
384  // see also program_manager.cpp::createSpirvProgram()
385  using ContextImplPtr = std::shared_ptr<sycl::detail::context_impl>;
387  ContextImplPtr ContextImpl = getSyclObjImpl(MContext);
388  const PluginPtr &Plugin = ContextImpl->getPlugin();
389  Plugin->call<PiApiKind::piProgramCreate>(
390  ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &PiProgram);
391  // program created by piProgramCreate is implicitly retained.
392 
393  std::vector<pi::PiDevice> DeviceVec;
394  DeviceVec.reserve(Devices.size());
395  for (const auto &SyclDev : Devices) {
396  pi::PiDevice Dev = getSyclObjImpl(SyclDev)->getHandleRef();
397  DeviceVec.push_back(Dev);
398  }
399  Plugin->call<errc::build, PiApiKind::piProgramBuild>(
400  PiProgram, DeviceVec.size(), DeviceVec.data(), nullptr, nullptr,
401  nullptr);
402 
403  // Get the number of kernels in the program.
404  size_t NumKernels;
405  Plugin->call<PiApiKind::piProgramGetInfo>(
406  PiProgram, PI_PROGRAM_INFO_NUM_KERNELS, sizeof(size_t), &NumKernels,
407  nullptr);
408 
409  // Get the kernel names.
410  size_t KernelNamesSize;
411  Plugin->call<PiApiKind::piProgramGetInfo>(
412  PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, 0, nullptr, &KernelNamesSize);
413 
414  // semi-colon delimited list of kernel names.
415  std::string KernelNamesStr(KernelNamesSize, ' ');
416  Plugin->call<PiApiKind::piProgramGetInfo>(
417  PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, KernelNamesStr.size(),
418  &KernelNamesStr[0], nullptr);
419  std::vector<std::string> KernelNames =
420  detail::split_string(KernelNamesStr, ';');
421 
422  // make the device image and the kernel_bundle_impl
423  auto KernelIDs = std::make_shared<std::vector<kernel_id>>();
424  auto DevImgImpl = std::make_shared<device_image_impl>(
425  nullptr, MContext, MDevices, bundle_state::executable, KernelIDs,
426  PiProgram);
427  device_image_plain DevImg{DevImgImpl};
428  return std::make_shared<kernel_bundle_impl>(MContext, MDevices, DevImg,
429  KernelNames);
430  }
431 
432  bool ext_oneapi_has_kernel(const std::string &Name) {
433  auto it = std::find(KernelNames.begin(), KernelNames.end(), Name);
434  return it != KernelNames.end();
435  }
436 
437  kernel
438  ext_oneapi_get_kernel(const std::string &Name,
439  const std::shared_ptr<kernel_bundle_impl> &Self) {
440  if (KernelNames.empty())
442  "'ext_oneapi_get_kernel' is only available in "
443  "kernel_bundles successfully built from "
444  "kernel_bundle<bundle_state:ext_oneapi_source>.");
445 
446  if (!ext_oneapi_has_kernel(Name))
448  "kernel '" + Name + "' not found in kernel_bundle");
449 
450  assert(MDeviceImages.size() > 0);
451  const std::shared_ptr<detail::device_image_impl> &DeviceImageImpl =
452  detail::getSyclObjImpl(MDeviceImages[0]);
453  sycl::detail::pi::PiProgram PiProgram = DeviceImageImpl->get_program_ref();
454  ContextImplPtr ContextImpl = getSyclObjImpl(MContext);
455  const PluginPtr &Plugin = ContextImpl->getPlugin();
457  Plugin->call<PiApiKind::piKernelCreate>(PiProgram, Name.c_str(), &PiKernel);
458  // Kernel created by piKernelCreate is implicitly retained.
459 
460  std::shared_ptr<kernel_impl> KernelImpl = std::make_shared<kernel_impl>(
461  PiKernel, detail::getSyclObjImpl(MContext), Self);
462 
463  return detail::createSyclObjFromImpl<kernel>(KernelImpl);
464  }
465 
466  bool empty() const noexcept { return MDeviceImages.empty(); }
467 
469  return MContext.get_platform().get_backend();
470  }
471 
472  context get_context() const noexcept { return MContext; }
473 
474  const std::vector<device> &get_devices() const noexcept { return MDevices; }
475 
476  std::vector<kernel_id> get_kernel_ids() const {
477  // Collect kernel ids from all device images, then remove duplicates
478 
479  std::vector<kernel_id> Result;
480  for (const device_image_plain &DeviceImage : MDeviceImages) {
481  const std::vector<kernel_id> &KernelIDs =
482  getSyclObjImpl(DeviceImage)->get_kernel_ids();
483 
484  Result.insert(Result.end(), KernelIDs.begin(), KernelIDs.end());
485  }
486  std::sort(Result.begin(), Result.end(), LessByNameComp{});
487 
488  auto NewIt = std::unique(Result.begin(), Result.end(), EqualByNameComp{});
489  Result.erase(NewIt, Result.end());
490 
491  return Result;
492  }
493 
494  kernel
495  get_kernel(const kernel_id &KernelID,
496  const std::shared_ptr<detail::kernel_bundle_impl> &Self) const {
497  using ImageImpl = std::shared_ptr<detail::device_image_impl>;
498  // Selected image.
499  ImageImpl SelectedImage = nullptr;
500  // Image where specialization constants are replaced with default values.
501  ImageImpl ImageWithReplacedSpecConsts = nullptr;
502  // Original image where specialization constants are not replaced with
503  // default values.
504  ImageImpl OriginalImage = nullptr;
505  // Used to track if any of the candidate images has specialization values
506  // set.
507  bool SpecConstsSet = false;
508  for (auto &DeviceImage : MDeviceImages) {
509  if (!DeviceImage.has_kernel(KernelID))
510  continue;
511 
512  const auto DeviceImageImpl = detail::getSyclObjImpl(DeviceImage);
513  SpecConstsSet |= DeviceImageImpl->is_any_specialization_constant_set();
514 
515  // Remember current image in corresponding variable depending on whether
516  // specialization constants are replaced with default value or not.
517  (DeviceImageImpl->specialization_constants_replaced_with_default()
518  ? ImageWithReplacedSpecConsts
519  : OriginalImage) = DeviceImageImpl;
520 
521  if (SpecConstsSet) {
522  // If specialization constant is set in any of the candidate images
523  // then we can't use ReplacedImage, so we select NativeImage if any or
524  // we select OriginalImage and keep iterating in case there is an image
525  // with native support.
526  SelectedImage = OriginalImage;
527  if (SelectedImage &&
528  SelectedImage->all_specialization_constant_native())
529  break;
530  } else {
531  // For now select ReplacedImage but it may be reset if any of the
532  // further device images has specialization constant value set. If after
533  // all iterations specialization constant values are not set in any of
534  // the candidate images then that will be the selected image.
535  // Also we don't want to use ReplacedImage if device image has native
536  // support.
537  if (ImageWithReplacedSpecConsts &&
538  !ImageWithReplacedSpecConsts->all_specialization_constant_native())
539  SelectedImage = ImageWithReplacedSpecConsts;
540  else
541  // In case if we don't have or don't use ReplacedImage.
542  SelectedImage = OriginalImage;
543  }
544  }
545 
546  if (!SelectedImage)
548  "The kernel bundle does not contain the kernel "
549  "identified by kernelId.");
550 
551  auto [Kernel, CacheMutex, ArgMask] =
553  MContext, KernelID.get_name(), /*PropList=*/{},
554  SelectedImage->get_program_ref());
555 
556  std::shared_ptr<kernel_impl> KernelImpl =
557  std::make_shared<kernel_impl>(Kernel, detail::getSyclObjImpl(MContext),
558  SelectedImage, Self, ArgMask, CacheMutex);
559 
560  return detail::createSyclObjFromImpl<kernel>(KernelImpl);
561  }
562 
563  bool has_kernel(const kernel_id &KernelID) const noexcept {
564  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
565  [&KernelID](const device_image_plain &DeviceImage) {
566  return DeviceImage.has_kernel(KernelID);
567  });
568  }
569 
570  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
571  return std::any_of(
572  MDeviceImages.begin(), MDeviceImages.end(),
573  [&KernelID, &Dev](const device_image_plain &DeviceImage) {
574  return DeviceImage.has_kernel(KernelID, Dev);
575  });
576  }
577 
579  return std::any_of(
580  MDeviceImages.begin(), MDeviceImages.end(),
581  [](const device_image_plain &DeviceImage) {
582  return getSyclObjImpl(DeviceImage)->has_specialization_constants();
583  });
584  }
585 
587  return contains_specialization_constants() &&
588  std::all_of(MDeviceImages.begin(), MDeviceImages.end(),
589  [](const device_image_plain &DeviceImage) {
590  return getSyclObjImpl(DeviceImage)
591  ->all_specialization_constant_native();
592  });
593  }
594 
595  bool has_specialization_constant(const char *SpecName) const noexcept {
596  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
597  [SpecName](const device_image_plain &DeviceImage) {
598  return getSyclObjImpl(DeviceImage)
599  ->has_specialization_constant(SpecName);
600  });
601  }
602 
603  void set_specialization_constant_raw_value(const char *SpecName,
604  const void *Value,
605  size_t Size) noexcept {
606  if (has_specialization_constant(SpecName))
607  for (const device_image_plain &DeviceImage : MDeviceImages)
608  getSyclObjImpl(DeviceImage)
609  ->set_specialization_constant_raw_value(SpecName, Value);
610  else {
611  std::vector<unsigned char> &Val = MSpecConstValues[std::string{SpecName}];
612  Val.resize(Size);
613  std::memcpy(Val.data(), Value, Size);
614  }
615  }
616 
617  void get_specialization_constant_raw_value(const char *SpecName,
618  void *ValueRet) const noexcept {
619  for (const device_image_plain &DeviceImage : MDeviceImages)
620  if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) {
621  getSyclObjImpl(DeviceImage)
622  ->get_specialization_constant_raw_value(SpecName, ValueRet);
623  return;
624  }
625 
626  // Specialization constant wasn't found in any of the device images,
627  // try to fetch value from kernel_bundle.
628  if (MSpecConstValues.count(std::string{SpecName}) != 0) {
629  const std::vector<unsigned char> &Val =
630  MSpecConstValues.at(std::string{SpecName});
631  auto *Dest = static_cast<unsigned char *>(ValueRet);
632  std::uninitialized_copy(Val.begin(), Val.end(), Dest);
633  return;
634  }
635 
636  assert(false &&
637  "get_specialization_constant_raw_value called for missing constant");
638  }
639 
640  bool is_specialization_constant_set(const char *SpecName) const noexcept {
641  bool SetInDevImg =
642  std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
643  [SpecName](const device_image_plain &DeviceImage) {
644  return getSyclObjImpl(DeviceImage)
645  ->is_specialization_constant_set(SpecName);
646  });
647  return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0;
648  }
649 
650  const device_image_plain *begin() const { return MDeviceImages.data(); }
651 
652  const device_image_plain *end() const {
653  return MDeviceImages.data() + MDeviceImages.size();
654  }
655 
656  size_t size() const noexcept { return MDeviceImages.size(); }
657 
658  bundle_state get_bundle_state() const { return MState; }
659 
660  const SpecConstMapT &get_spec_const_map_ref() const noexcept {
661  return MSpecConstValues;
662  }
663 
664  bool isInterop() const { return MIsInterop; }
665 
666  bool add_kernel(const kernel_id &KernelID, const device &Dev) {
667  // Skip if kernel is already there
668  if (has_kernel(KernelID, Dev))
669  return true;
670 
671  // First try and get images in current bundle state
672  const bundle_state BundleState = get_bundle_state();
673  std::vector<device_image_plain> NewDevImgs =
675  MContext, {Dev}, {KernelID}, BundleState);
676 
677  // No images found so we report as not inserted
678  if (NewDevImgs.empty())
679  return false;
680 
681  // Propagate already set specialization constants to the new images
682  for (device_image_plain &DevImg : NewDevImgs)
683  for (auto SpecConst : MSpecConstValues)
684  getSyclObjImpl(DevImg)->set_specialization_constant_raw_value(
685  SpecConst.first.c_str(), SpecConst.second.data());
686 
687  // Add the images to the collection
688  MDeviceImages.insert(MDeviceImages.end(), NewDevImgs.begin(),
689  NewDevImgs.end());
690  return true;
691  }
692 
693 private:
694  context MContext;
695  std::vector<device> MDevices;
696  std::vector<device_image_plain> MDeviceImages;
697  // This map stores values for specialization constants, that are missing
698  // from any device image.
699  SpecConstMapT MSpecConstValues;
700  bool MIsInterop = false;
701  bundle_state MState;
702  // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames
703  const syclex::source_language Language = syclex::source_language::opencl;
704  const std::variant<std::string, std::vector<std::byte>> Source;
705  // only kernel_bundles created from source have KernelNames member.
706  std::vector<std::string> KernelNames;
707 };
708 
709 } // namespace detail
710 } // namespace _V1
711 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:51
std::vector< device_image_plain > getSYCLDeviceImages(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
static ProgramManager & getInstance()
std::tuple< sycl::detail::pi::PiKernel, std::mutex *, const KernelArgMask *, sycl::detail::pi::PiProgram > getOrCreateKernel(const ContextImplPtr &ContextImpl, const DeviceImplPtr &DeviceImpl, const std::string &KernelName, const NDRDescT &NDRDesc={})
std::vector< device_image_plain > link(const device_image_plain &DeviceImages, const std::vector< device > &Devs, const property_list &PropList)
The class is an impl counterpart of the sycl::kernel_bundle.
bool has_kernel(const kernel_id &KernelID) const noexcept
const SpecConstMapT & get_spec_const_map_ref() const noexcept
std::shared_ptr< kernel_bundle_impl > build_from_source(const std::vector< device > Devices, const std::vector< std::string > &BuildOptions, std::string *LogPtr)
kernel_bundle_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, std::vector< device > Devs, const property_list &PropList)
kernel_bundle_impl(context Ctx, std::vector< device > Devs)
bool add_kernel(const kernel_id &KernelID, const device &Dev)
kernel ext_oneapi_get_kernel(const std::string &Name, const std::shared_ptr< kernel_bundle_impl > &Self)
kernel_bundle_impl(context Ctx, std::vector< device > Devs, device_image_plain &DevImage)
kernel_bundle_impl(const context &Context, syclex::source_language Lang, const std::string &Src)
const device_image_plain * end() const
kernel_bundle_impl(const context &Context, syclex::source_language Lang, const std::vector< std::byte > &Bytes)
const std::vector< device > & get_devices() const noexcept
void set_specialization_constant_raw_value(const char *SpecName, const void *Value, size_t Size) noexcept
bool contains_specialization_constants() const noexcept
std::vector< kernel_id > get_kernel_ids() const
bool is_specialization_constant_set(const char *SpecName) const noexcept
bool native_specialization_constant() const noexcept
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const std::vector< kernel_id > &KernelIDs, bundle_state State)
kernel_bundle_impl(const kernel_bundle< bundle_state::input > &InputBundle, std::vector< device > Devs, const property_list &PropList, bundle_state TargetState)
const device_image_plain * begin() const
kernel_bundle_impl(context Ctx, std::vector< device > Devs, bundle_state State)
void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet) const noexcept
bool has_specialization_constant(const char *SpecName) const noexcept
bool ext_oneapi_has_kernel(const std::string &Name)
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
kernel_bundle_impl(context Ctx, std::vector< device > Devs, device_image_plain &DevImage, std::vector< std::string > KNames)
kernel get_kernel(const kernel_id &KernelID, const std::shared_ptr< detail::kernel_bundle_impl > &Self) const
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const DevImgSelectorImpl &Selector, bundle_state State)
kernel_bundle_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:59
The kernel_bundle class represents collection of device images in a particular state.
Objects of the class identify kernel is some kernel_bundle related APIs.
const char * get_name() const noexcept
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:74
Objects of the property_list class are containers for the SYCL properties.
::pi_kernel PiKernel
Definition: pi.hpp:138
::pi_program PiProgram
Definition: pi.hpp:137
std::shared_ptr< device_image_impl > DeviceImageImplPtr
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
Definition: event_impl.hpp:33
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
std::vector< std::string > split_string(const std::string &str, char delimeter)
Definition: common.cpp:74
std::shared_ptr< plugin > PluginPtr
Definition: pi.hpp:48
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
Function for_each(Group g, Ptr first, Ptr last, Function f)
spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector< std::string > &UserArgs, std::string *LogPtr)
kernel_bundle< bundle_state::executable > build(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:94
Definition: access.hpp:18
pi_result piKernelCreate(pi_program program, const char *kernel_name, pi_kernel *ret_kernel)
Definition: pi_cuda.cpp:341
pi_result piProgramGetInfo(pi_program program, pi_program_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_cuda.cpp:272
pi_result piProgramBuild(pi_program program, pi_uint32 num_devices, const pi_device *device_list, const char *options, void(*pfn_notify)(pi_program program, void *user_data), void *user_data)
pi_result piProgramCreate(pi_context context, const void *il, size_t length, pi_program *res_program)
Definition: pi_cuda.cpp:248
@ PI_PROGRAM_INFO_KERNEL_NAMES
Definition: pi.h:455
@ PI_PROGRAM_INFO_NUM_KERNELS
Definition: pi.h:454
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept