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 
12 #include <CL/sycl/context.hpp>
14 #include <CL/sycl/detail/pi.h>
15 #include <CL/sycl/device.hpp>
18 #include <detail/kernel_impl.hpp>
20 
21 #include <algorithm>
22 #include <cassert>
23 #include <cstring>
24 #include <memory>
25 #include <vector>
26 
28 namespace sycl {
29 namespace detail {
30 
31 static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,
32  const context &Context) {
33  const std::vector<device> &ContextDevices = Context.get_devices();
34  return std::all_of(
35  Devices.begin(), Devices.end(), [&ContextDevices](const device &Dev) {
36  return ContextDevices.end() !=
37  std::find(ContextDevices.begin(), ContextDevices.end(), Dev);
38  });
39 }
40 
41 static bool checkAllDevicesHaveAspect(const std::vector<device> &Devices,
42  aspect Aspect) {
43  return std::all_of(Devices.begin(), Devices.end(),
44  [&Aspect](const device &Dev) { return Dev.has(Aspect); });
45 }
46 
47 // The class is an impl counterpart of the sycl::kernel_bundle.
48 // It provides an access and utilities to manage set of sycl::device_images
49 // objects.
51 
52  using SpecConstMapT = std::map<std::string, std::vector<unsigned char>>;
53 
54  void common_ctor_checks(bundle_state State) {
55  const bool AllDevicesInTheContext =
56  checkAllDevicesAreInContext(MDevices, MContext);
57  if (MDevices.empty() || !AllDevicesInTheContext)
58  throw sycl::exception(
59  make_error_code(errc::invalid),
60  "Not all devices are associated with the context or "
61  "vector of devices is empty");
62 
63  if (bundle_state::input == State &&
64  !checkAllDevicesHaveAspect(MDevices, aspect::online_compiler))
65  throw sycl::exception(make_error_code(errc::invalid),
66  "Not all devices have aspect::online_compiler");
67 
68  if (bundle_state::object == State &&
69  !checkAllDevicesHaveAspect(MDevices, aspect::online_linker))
70  throw sycl::exception(make_error_code(errc::invalid),
71  "Not all devices have aspect::online_linker");
72  }
73 
74 public:
75  kernel_bundle_impl(context Ctx, std::vector<device> Devs, bundle_state State)
76  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
77 
78  common_ctor_checks(State);
79 
80  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
81  MContext, MDevices, State);
82  }
83 
84  // Interop constructor used by make_kernel
85  kernel_bundle_impl(context Ctx, std::vector<device> Devs)
86  : MContext(Ctx), MDevices(Devs), MState(bundle_state::executable) {
87  if (!checkAllDevicesAreInContext(Devs, Ctx))
88  throw sycl::exception(
89  make_error_code(errc::invalid),
90  "Not all devices are associated with the context or "
91  "vector of devices is empty");
92  MIsInterop = true;
93  }
94 
95  // Interop constructor
96  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
97  device_image_plain &DevImage)
98  : kernel_bundle_impl(Ctx, Devs) {
99  MDeviceImages.push_back(DevImage);
100  }
101 
102  // Matches sycl::build and sycl::compile
103  // Have one constructor because sycl::build and sycl::compile have the same
104  // signature
106  std::vector<device> Devs, const property_list &PropList,
107  bundle_state TargetState)
108  : MContext(InputBundle.get_context()), MDevices(std::move(Devs)),
109  MState(TargetState) {
110 
111  MSpecConstValues = getSyclObjImpl(InputBundle)->get_spec_const_map_ref();
112 
113  const std::vector<device> &InputBundleDevices =
114  getSyclObjImpl(InputBundle)->get_devices();
115  const bool AllDevsAssociatedWithInputBundle =
116  std::all_of(MDevices.begin(), MDevices.end(),
117  [&InputBundleDevices](const device &Dev) {
118  return InputBundleDevices.end() !=
119  std::find(InputBundleDevices.begin(),
120  InputBundleDevices.end(), Dev);
121  });
122  if (MDevices.empty() || !AllDevsAssociatedWithInputBundle)
123  throw sycl::exception(
124  make_error_code(errc::invalid),
125  "Not all devices are in the set of associated "
126  "devices for input bundle or vector of devices is empty");
127 
128  for (const device_image_plain &DeviceImage : InputBundle) {
129  // Skip images which are not compatible with devices provided
130  if (std::none_of(
131  MDevices.begin(), MDevices.end(),
132  [&DeviceImage](const device &Dev) {
133  return getSyclObjImpl(DeviceImage)->compatible_with_device(Dev);
134  }))
135  continue;
136 
137  switch (TargetState) {
138  case bundle_state::object:
139  MDeviceImages.push_back(detail::ProgramManager::getInstance().compile(
140  DeviceImage, MDevices, PropList));
141  break;
142  case bundle_state::executable:
143  MDeviceImages.push_back(detail::ProgramManager::getInstance().build(
144  DeviceImage, MDevices, PropList));
145  break;
146  case bundle_state::input:
147  throw sycl::runtime_error(
148  "Internal error. The target state should not be input",
150  break;
151  }
152  }
153  }
154 
155  // Matches sycl::link
157  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
158  std::vector<device> Devs, const property_list &PropList)
159  : MDevices(std::move(Devs)), MState(bundle_state::executable) {
160 
161  if (MDevices.empty())
162  throw sycl::exception(make_error_code(errc::invalid),
163  "Vector of devices is empty");
164 
165  if (ObjectBundles.empty())
166  return;
167 
168  MContext = ObjectBundles[0].get_context();
169  for (size_t I = 1; I < ObjectBundles.size(); ++I) {
170  if (ObjectBundles[I].get_context() != MContext)
171  throw sycl::exception(
172  make_error_code(errc::invalid),
173  "Not all input bundles have the same associated context");
174  }
175 
176  // Check if any of the devices in devs are not in the set of associated
177  // devices for any of the bundles in ObjectBundles
178  const bool AllDevsAssociatedWithInputBundles = std::all_of(
179  MDevices.begin(), MDevices.end(), [&ObjectBundles](const device &Dev) {
180  // Number of devices is expected to be small
181  return std::all_of(
182  ObjectBundles.begin(), ObjectBundles.end(),
183  [&Dev](const kernel_bundle<bundle_state::object> &KernelBundle) {
184  const std::vector<device> &BundleDevices =
185  getSyclObjImpl(KernelBundle)->get_devices();
186  return BundleDevices.end() != std::find(BundleDevices.begin(),
187  BundleDevices.end(),
188  Dev);
189  });
190  });
191  if (!AllDevsAssociatedWithInputBundles)
192  throw sycl::exception(make_error_code(errc::invalid),
193  "Not all devices are in the set of associated "
194  "devices for input bundles");
195 
196  // TODO: Unify with c'tor for sycl::comile and sycl::build by calling
197  // sycl::join on vector of kernel_bundles
198 
199  // The loop below just links each device image separately, not linking any
200  // two device images together. This is correct so long as each device image
201  // has no unresolved symbols. That's the case when device images are created
202  // from generic SYCL APIs. There's no way in generic SYCL to create a kernel
203  // which references an undefined symbol. If we decide in the future to allow
204  // a backend interop API to create a "sycl::kernel_bundle" that references
205  // undefined symbols, then the logic in this loop will need to be changed.
206  for (const kernel_bundle<bundle_state::object> &ObjectBundle :
207  ObjectBundles) {
208  for (const device_image_plain &DeviceImage : ObjectBundle) {
209 
210  // Skip images which are not compatible with devices provided
211  if (std::none_of(MDevices.begin(), MDevices.end(),
212  [&DeviceImage](const device &Dev) {
213  return getSyclObjImpl(DeviceImage)
214  ->compatible_with_device(Dev);
215  }))
216  continue;
217 
218  const std::vector<device_image_plain> VectorOfOneImage{DeviceImage};
219  std::vector<device_image_plain> LinkedResults =
220  detail::ProgramManager::getInstance().link(VectorOfOneImage,
221  MDevices, PropList);
222  MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
223  LinkedResults.end());
224  }
225  }
226 
227  for (const kernel_bundle<bundle_state::object> &Bundle : ObjectBundles) {
228  const KernelBundleImplPtr BundlePtr = getSyclObjImpl(Bundle);
229  for (const std::pair<const std::string, std::vector<unsigned char>>
230  &SpecConst : BundlePtr->MSpecConstValues) {
231  MSpecConstValues[SpecConst.first] = SpecConst.second;
232  }
233  }
234  }
235 
236  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
237  const std::vector<kernel_id> &KernelIDs,
238  bundle_state State)
239  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
240 
241  // TODO: Add a check that all kernel ids are compatible with at least one
242  // device in Devs
243  common_ctor_checks(State);
244 
245  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
246  MContext, MDevices, KernelIDs, State);
247  }
248 
249  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
250  const DevImgSelectorImpl &Selector, bundle_state State)
251  : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
252 
253  common_ctor_checks(State);
254 
255  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
256  MContext, MDevices, Selector, State);
257  }
258 
259  // C'tor matches sycl::join API
260  kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
261  bundle_state State)
262  : MState(State) {
263  if (Bundles.empty())
264  return;
265 
266  MContext = Bundles[0]->MContext;
267  MDevices = Bundles[0]->MDevices;
268  for (size_t I = 1; I < Bundles.size(); ++I) {
269  if (Bundles[I]->MContext != MContext)
270  throw sycl::exception(
271  make_error_code(errc::invalid),
272  "Not all input bundles have the same associated context.");
273  if (Bundles[I]->MDevices != MDevices)
274  throw sycl::exception(
275  make_error_code(errc::invalid),
276  "Not all input bundles have the same set of associated devices.");
277  }
278 
279  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
280 
281  MDeviceImages.insert(MDeviceImages.end(), Bundle->MDeviceImages.begin(),
282  Bundle->MDeviceImages.end());
283  }
284 
285  std::sort(MDeviceImages.begin(), MDeviceImages.end(),
287 
288  if (get_bundle_state() == bundle_state::input) {
289  // Copy spec constants values from the device images to be removed.
290  auto MergeSpecConstants = [this](const device_image_plain &Img) {
291  const detail::DeviceImageImplPtr &ImgImpl = getSyclObjImpl(Img);
292  const std::map<std::string,
293  std::vector<device_image_impl::SpecConstDescT>>
294  &SpecConsts = ImgImpl->get_spec_const_data_ref();
295  const std::vector<unsigned char> &Blob =
296  ImgImpl->get_spec_const_blob_ref();
297  for (const std::pair<const std::string,
298  std::vector<device_image_impl::SpecConstDescT>>
299  &SpecConst : SpecConsts) {
300  if (SpecConst.second.front().IsSet)
301  set_specialization_constant_raw_value(
302  SpecConst.first.c_str(),
303  Blob.data() + SpecConst.second.front().BlobOffset,
304  SpecConst.second.back().CompositeOffset +
305  SpecConst.second.back().Size);
306  }
307  };
308  std::for_each(MDeviceImages.begin(), MDeviceImages.end(),
309  MergeSpecConstants);
310  }
311 
312  const auto DevImgIt =
313  std::unique(MDeviceImages.begin(), MDeviceImages.end());
314 
315  // Remove duplicate device images.
316  MDeviceImages.erase(DevImgIt, MDeviceImages.end());
317 
318  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
319  for (const std::pair<const std::string, std::vector<unsigned char>>
320  &SpecConst : Bundle->MSpecConstValues) {
321  set_specialization_constant_raw_value(SpecConst.first.c_str(),
322  SpecConst.second.data(),
323  SpecConst.second.size());
324  }
325  }
326  }
327 
328  bool empty() const noexcept { return MDeviceImages.empty(); }
329 
330  backend get_backend() const noexcept {
331  return MContext.get_platform().get_backend();
332  }
333 
334  context get_context() const noexcept { return MContext; }
335 
336  const std::vector<device> &get_devices() const noexcept { return MDevices; }
337 
338  std::vector<kernel_id> get_kernel_ids() const {
339  // Collect kernel ids from all device images, then remove duplicates
340 
341  std::vector<kernel_id> Result;
342  for (const device_image_plain &DeviceImage : MDeviceImages) {
343  const std::vector<kernel_id> &KernelIDs =
344  getSyclObjImpl(DeviceImage)->get_kernel_ids();
345 
346  Result.insert(Result.end(), KernelIDs.begin(), KernelIDs.end());
347  }
348  std::sort(Result.begin(), Result.end(), LessByNameComp{});
349 
350  auto NewIt = std::unique(Result.begin(), Result.end(), EqualByNameComp{});
351  Result.erase(NewIt, Result.end());
352 
353  return Result;
354  }
355 
356  kernel
357  get_kernel(const kernel_id &KernelID,
358  const std::shared_ptr<detail::kernel_bundle_impl> &Self) const {
359 
360  auto It = std::find_if(MDeviceImages.begin(), MDeviceImages.end(),
361  [&KernelID](const device_image_plain &DeviceImage) {
362  return DeviceImage.has_kernel(KernelID);
363  });
364 
365  if (MDeviceImages.end() == It)
366  throw sycl::exception(make_error_code(errc::invalid),
367  "The kernel bundle does not contain the kernel "
368  "identified by kernelId.");
369 
370  const std::shared_ptr<detail::device_image_impl> &DeviceImageImpl =
372 
373  RT::PiKernel Kernel = nullptr;
374  std::tie(Kernel, std::ignore) =
375  detail::ProgramManager::getInstance().getOrCreateKernel(
376  MContext, KernelID.get_name(), /*PropList=*/{},
377  DeviceImageImpl->get_program_ref());
378 
379  std::shared_ptr<kernel_impl> KernelImpl = std::make_shared<kernel_impl>(
380  Kernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self);
381 
382  return detail::createSyclObjFromImpl<kernel>(KernelImpl);
383  }
384 
385  bool has_kernel(const kernel_id &KernelID) const noexcept {
386  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
387  [&KernelID](const device_image_plain &DeviceImage) {
388  return DeviceImage.has_kernel(KernelID);
389  });
390  }
391 
392  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
393  return std::any_of(
394  MDeviceImages.begin(), MDeviceImages.end(),
395  [&KernelID, &Dev](const device_image_plain &DeviceImage) {
396  return DeviceImage.has_kernel(KernelID, Dev);
397  });
398  }
399 
400  bool contains_specialization_constants() const noexcept {
401  return std::any_of(
402  MDeviceImages.begin(), MDeviceImages.end(),
403  [](const device_image_plain &DeviceImage) {
404  return getSyclObjImpl(DeviceImage)->has_specialization_constants();
405  });
406  }
407 
408  bool native_specialization_constant() const noexcept {
409  return std::all_of(MDeviceImages.begin(), MDeviceImages.end(),
410  [](const device_image_plain &DeviceImage) {
411  return getSyclObjImpl(DeviceImage)
412  ->all_specialization_constant_native();
413  });
414  }
415 
416  bool has_specialization_constant(const char *SpecName) const noexcept {
417  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
418  [SpecName](const device_image_plain &DeviceImage) {
419  return getSyclObjImpl(DeviceImage)
420  ->has_specialization_constant(SpecName);
421  });
422  }
423 
424  void set_specialization_constant_raw_value(const char *SpecName,
425  const void *Value,
426  size_t Size) noexcept {
427  if (has_specialization_constant(SpecName))
428  for (const device_image_plain &DeviceImage : MDeviceImages)
429  getSyclObjImpl(DeviceImage)
430  ->set_specialization_constant_raw_value(SpecName, Value);
431  else {
432  const auto *DataPtr = static_cast<const unsigned char *>(Value);
433  std::vector<unsigned char> &Val = MSpecConstValues[std::string{SpecName}];
434  Val.resize(Size);
435  Val.insert(Val.begin(), DataPtr, DataPtr + Size);
436  }
437  }
438 
439  void get_specialization_constant_raw_value(const char *SpecName,
440  void *ValueRet) const noexcept {
441  for (const device_image_plain &DeviceImage : MDeviceImages)
442  if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) {
443  getSyclObjImpl(DeviceImage)
444  ->get_specialization_constant_raw_value(SpecName, ValueRet);
445  return;
446  }
447 
448  // Specialization constant wasn't found in any of the device images,
449  // try to fetch value from kernel_bundle.
450  if (MSpecConstValues.count(std::string{SpecName}) != 0) {
451  const std::vector<unsigned char> &Val =
452  MSpecConstValues.at(std::string{SpecName});
453  auto *Dest = static_cast<unsigned char *>(ValueRet);
454  std::uninitialized_copy(Val.begin(), Val.end(), Dest);
455  return;
456  }
457 
458  assert(false &&
459  "get_specialization_constant_raw_value called for missing constant");
460  }
461 
462  bool is_specialization_constant_set(const char *SpecName) const noexcept {
463  bool SetInDevImg =
464  std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
465  [SpecName](const device_image_plain &DeviceImage) {
466  return getSyclObjImpl(DeviceImage)
467  ->is_specialization_constant_set(SpecName);
468  });
469  return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0;
470  }
471 
472  const device_image_plain *begin() const { return MDeviceImages.data(); }
473 
474  const device_image_plain *end() const {
475  return MDeviceImages.data() + MDeviceImages.size();
476  }
477 
478  size_t size() const noexcept { return MDeviceImages.size(); }
479 
480  bundle_state get_bundle_state() const { return MState; }
481 
482  const SpecConstMapT &get_spec_const_map_ref() const noexcept {
483  return MSpecConstValues;
484  }
485 
486  bool isInterop() const { return MIsInterop; }
487 
488  bool add_kernel(const kernel_id &KernelID, const device &Dev) {
489  // Skip if kernel is already there
490  if (has_kernel(KernelID, Dev))
491  return true;
492 
493  // First try and get images in current bundle state
494  const bundle_state BundleState = get_bundle_state();
495  std::vector<device_image_plain> NewDevImgs =
496  detail::ProgramManager::getInstance().getSYCLDeviceImages(
497  MContext, {Dev}, {KernelID}, BundleState);
498 
499  // No images found so we report as not inserted
500  if (NewDevImgs.empty())
501  return false;
502 
503  // Propagate already set specialization constants to the new images
504  for (device_image_plain &DevImg : NewDevImgs)
505  for (auto SpecConst : MSpecConstValues)
506  getSyclObjImpl(DevImg)->set_specialization_constant_raw_value(
507  SpecConst.first.c_str(), SpecConst.second.data());
508 
509  // Add the images to the collection
510  MDeviceImages.insert(MDeviceImages.end(), NewDevImgs.begin(),
511  NewDevImgs.end());
512  return true;
513  }
514 
515 private:
516  context MContext;
517  std::vector<device> MDevices;
518  std::vector<device_image_plain> MDeviceImages;
519  // This map stores values for specialization constants, that are missing
520  // from any device image.
521  SpecConstMapT MSpecConstValues;
522  bool MIsInterop = false;
523  bundle_state MState;
524 };
525 
526 } // namespace detail
527 } // namespace sycl
528 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::backend
backend
Definition: backend_types.hpp:21
pi.h
cl::sycl::detail::kernel_bundle_impl::size
size_t size() const noexcept
Definition: kernel_bundle_impl.hpp:478
cl::sycl::kernel_bundle
The kernel_bundle class represents collection of device images in a particular state.
Definition: kernel.hpp:28
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const DevImgSelectorImpl &Selector, bundle_state State)
Definition: kernel_bundle_impl.hpp:249
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(context Ctx, std::vector< device > Devs, bundle_state State)
Definition: kernel_bundle_impl.hpp:75
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:88
cl::sycl::detail::kernel_bundle_impl
Definition: kernel_bundle_impl.hpp:50
cl::sycl::build
kernel_bundle< bundle_state::executable > build(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
Definition: kernel_bundle.hpp:714
cl::sycl::detail::kernel_bundle_impl::add_kernel
bool add_kernel(const kernel_id &KernelID, const device &Dev)
Definition: kernel_bundle_impl.hpp:488
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(context Ctx, std::vector< device > Devs, const std::vector< kernel_id > &KernelIDs, bundle_state State)
Definition: kernel_bundle_impl.hpp:236
cl::sycl::detail::for_each
Function for_each(Group g, Ptr first, Ptr last, Function f)
Definition: group_algorithm.hpp:148
device.hpp
cl::sycl::detail::device_image_plain
Definition: kernel_bundle.hpp:71
cl::sycl::detail::kernel_bundle_impl::empty
bool empty() const noexcept
Definition: kernel_bundle_impl.hpp:328
cl::sycl::detail::kernel_bundle_impl::get_spec_const_map_ref
const SpecConstMapT & get_spec_const_map_ref() const noexcept
Definition: kernel_bundle_impl.hpp:482
cl::sycl::detail::kernel_bundle_impl::end
const device_image_plain * end() const
Definition: kernel_bundle_impl.hpp:474
context.hpp
cl::sycl::make_error_code
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:121
cl::sycl::detail::KernelBundleImplPtr
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
Definition: kernel_bundle.hpp:132
sycl
Definition: invoke_simd.hpp:68
cl::sycl::detail::kernel_bundle_impl::has_specialization_constant
bool has_specialization_constant(const char *SpecName) const noexcept
Definition: kernel_bundle_impl.hpp:416
cl::sycl::detail::checkAllDevicesHaveAspect
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
Definition: kernel_bundle_impl.hpp:41
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:26
cl::sycl::detail::kernel_bundle_impl::isInterop
bool isInterop() const
Definition: kernel_bundle_impl.hpp:486
cl::sycl::detail::kernel_bundle_impl::get_kernel
kernel get_kernel(const kernel_id &KernelID, const std::shared_ptr< detail::kernel_bundle_impl > &Self) const
Definition: kernel_bundle_impl.hpp:357
cl::sycl::detail::kernel_bundle_impl::get_devices
const std::vector< device > & get_devices() const noexcept
Definition: kernel_bundle_impl.hpp:336
cl::sycl::bundle_state
bundle_state
Definition: kernel_bundle_enums.hpp:14
_pi_kernel
Implementation of a PI Kernel for CUDA.
Definition: pi_cuda.hpp:626
cl::sycl::detail::kernel_bundle_impl::get_context
context get_context() const noexcept
Definition: kernel_bundle_impl.hpp:334
cl::sycl::detail::checkAllDevicesAreInContext
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
Definition: kernel_bundle_impl.hpp:31
cl::sycl::detail::kernel_bundle_impl::is_specialization_constant_set
bool is_specialization_constant_set(const char *SpecName) const noexcept
Definition: kernel_bundle_impl.hpp:462
cl::sycl::detail::kernel_bundle_impl::begin
const device_image_plain * begin() const
Definition: kernel_bundle_impl.hpp:472
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:35
cl::sycl::detail::kernel_bundle_impl::has_kernel
bool has_kernel(const kernel_id &KernelID) const noexcept
Definition: kernel_bundle_impl.hpp:385
kernel_bundle.hpp
cl::sycl::compile
kernel_bundle< bundle_state::object > compile(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList={})
Definition: kernel_bundle.hpp:632
cl::sycl::detail::kernel_bundle_impl::has_kernel
bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept
Definition: kernel_bundle_impl.hpp:392
cl::sycl::kernel_id
Objects of the class identify kernel is some kernel_bundle related APIs.
Definition: kernel_bundle.hpp:39
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
Definition: kernel_bundle_impl.hpp:260
cl::sycl::bundle_state::executable
@ executable
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::kernel_id::get_name
const char * get_name() const noexcept
Definition: kernel_bundle.cpp:20
cl::sycl::detail::kernel_bundle_impl::get_specialization_constant_raw_value
void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet) const noexcept
Definition: kernel_bundle_impl.hpp:439
cl::sycl::detail::tie
auto tie(Ts &... Args)
Definition: tuple.hpp:40
none_of
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept
cl::sycl::aspect
aspect
Definition: aspects.hpp:15
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(context Ctx, std::vector< device > Devs, device_image_plain &DevImage)
Definition: kernel_bundle_impl.hpp:96
program_manager.hpp
cl::sycl::detail::kernel_bundle_impl::get_backend
backend get_backend() const noexcept
Definition: kernel_bundle_impl.hpp:330
cl::sycl::context::get_devices
std::vector< device > get_devices() const
Gets devices associated with this SYCL context.
Definition: context.cpp:127
cl::sycl::detail::LessByHash
Definition: device_image_impl.hpp:35
cl::sycl::detail::kernel_bundle_impl::get_bundle_state
bundle_state get_bundle_state() const
Definition: kernel_bundle_impl.hpp:480
all_of
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
cl::sycl::detail::EqualByNameComp
Definition: kernel_id_impl.hpp:23
device_image_impl.hpp
cl::sycl::detail::DevImgSelectorImpl
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
Definition: kernel_bundle.hpp:467
cl::sycl::detail::kernel_bundle_impl::native_specialization_constant
bool native_specialization_constant() const noexcept
Definition: kernel_bundle_impl.hpp:408
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:204
backend_types.hpp
std
Definition: accessor.hpp:2616
cl::sycl::detail::kernel_bundle_impl::set_specialization_constant_raw_value
void set_specialization_constant_raw_value(const char *SpecName, const void *Value, size_t Size) noexcept
Definition: kernel_bundle_impl.hpp:424
cl::sycl::detail::kernel_bundle_impl::get_kernel_ids
std::vector< kernel_id > get_kernel_ids() const
Definition: kernel_bundle_impl.hpp:338
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(context Ctx, std::vector< device > Devs)
Definition: kernel_bundle_impl.hpp:85
cl::sycl::context
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:35
any_of
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
common.hpp
cl::sycl::exception
Definition: exception.hpp:63
kernel_impl.hpp
cl::sycl::kernel
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:67
cl::sycl::detail::LessByNameComp
Definition: kernel_id_impl.hpp:16
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, std::vector< device > Devs, const property_list &PropList)
Definition: kernel_bundle_impl.hpp:156
cl::sycl::detail::kernel_bundle_impl::contains_specialization_constants
bool contains_specialization_constants() const noexcept
Definition: kernel_bundle_impl.hpp:400
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(const kernel_bundle< bundle_state::input > &InputBundle, std::vector< device > Devs, const property_list &PropList, bundle_state TargetState)
Definition: kernel_bundle_impl.hpp:105
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::detail::DeviceImageImplPtr
std::shared_ptr< device_image_impl > DeviceImageImplPtr
Definition: kernel_bundle.hpp:67