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