DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
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 template <class T> struct LessByHash {
32  bool operator()(const T &LHS, const T &RHS) const {
33  return getSyclObjImpl(LHS) < getSyclObjImpl(RHS);
34  }
35 };
36 
37 static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,
38  const context &Context) {
39  const std::vector<device> &ContextDevices = Context.get_devices();
40  return std::all_of(
41  Devices.begin(), Devices.end(), [&ContextDevices](const device &Dev) {
42  return ContextDevices.end() !=
43  std::find(ContextDevices.begin(), ContextDevices.end(), Dev);
44  });
45 }
46 
47 static bool checkAllDevicesHaveAspect(const std::vector<device> &Devices,
48  aspect Aspect) {
49  return std::all_of(Devices.begin(), Devices.end(),
50  [&Aspect](const device &Dev) { return Dev.has(Aspect); });
51 }
52 
53 // The class is an impl counterpart of the sycl::kernel_bundle.
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(
65  make_error_code(errc::invalid),
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))
71  throw sycl::exception(make_error_code(errc::invalid),
72  "Not all devices have aspect::online_compiler");
73 
74  if (bundle_state::object == State &&
75  !checkAllDevicesHaveAspect(MDevices, aspect::online_linker))
76  throw sycl::exception(make_error_code(errc::invalid),
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)) {
83 
84  common_ctor_checks(State);
85 
86  MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(
87  MContext, MDevices, State);
88  }
89 
90  // Interop constructor
91  kernel_bundle_impl(context Ctx, std::vector<device> Devs,
92  device_image_plain &DevImage)
93  : MContext(Ctx), MDevices(Devs) {
94  if (!checkAllDevicesAreInContext(Devs, Ctx))
95  throw sycl::exception(
96  make_error_code(errc::invalid),
97  "Not all devices are associated with the context or "
98  "vector of devices is empty");
99  MDeviceImages.push_back(DevImage);
100  MIsInterop = true;
101  }
102 
103  // Matches sycl::build and sycl::compile
104  // Have one constructor because sycl::build and sycl::compile have the same
105  // signature
107  std::vector<device> Devs, const property_list &PropList,
108  bundle_state TargetState)
109  : MContext(InputBundle.get_context()), MDevices(std::move(Devs)) {
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)) {
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)) {
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)) {
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  if (Bundles.empty())
262  return;
263 
264  MContext = Bundles[0]->MContext;
265  MDevices = Bundles[0]->MDevices;
266  for (size_t I = 1; I < Bundles.size(); ++I) {
267  if (Bundles[I]->MContext != MContext)
268  throw sycl::exception(
269  make_error_code(errc::invalid),
270  "Not all input bundles have the same associated context.");
271  if (Bundles[I]->MDevices != MDevices)
272  throw sycl::exception(
273  make_error_code(errc::invalid),
274  "Not all input bundles have the same set of associated devices.");
275  }
276 
277  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
278 
279  MDeviceImages.insert(MDeviceImages.end(), Bundle->MDeviceImages.begin(),
280  Bundle->MDeviceImages.end());
281  }
282 
283  std::sort(MDeviceImages.begin(), MDeviceImages.end(),
285 
286  if (get_bundle_state() == bundle_state::input) {
287  // Copy spec constants values from the device images to be removed.
288  auto MergeSpecConstants = [this](const device_image_plain &Img) {
289  const detail::DeviceImageImplPtr &ImgImpl = getSyclObjImpl(Img);
290  const std::map<std::string,
291  std::vector<device_image_impl::SpecConstDescT>>
292  &SpecConsts = ImgImpl->get_spec_const_data_ref();
293  const std::vector<unsigned char> &Blob =
294  ImgImpl->get_spec_const_blob_ref();
295  for (const std::pair<const std::string,
296  std::vector<device_image_impl::SpecConstDescT>>
297  &SpecConst : SpecConsts) {
298  if (SpecConst.second.front().IsSet)
299  set_specialization_constant_raw_value(
300  SpecConst.first.c_str(),
301  Blob.data() + SpecConst.second.front().BlobOffset,
302  SpecConst.second.back().CompositeOffset +
303  SpecConst.second.back().Size);
304  }
305  };
306  std::for_each(MDeviceImages.begin(), MDeviceImages.end(),
307  MergeSpecConstants);
308  }
309 
310  const auto DevImgIt =
311  std::unique(MDeviceImages.begin(), MDeviceImages.end());
312 
313  // Remove duplicate device images.
314  MDeviceImages.erase(DevImgIt, MDeviceImages.end());
315 
316  for (const detail::KernelBundleImplPtr &Bundle : Bundles) {
317  for (const std::pair<const std::string, std::vector<unsigned char>>
318  &SpecConst : Bundle->MSpecConstValues) {
319  set_specialization_constant_raw_value(SpecConst.first.c_str(),
320  SpecConst.second.data(),
321  SpecConst.second.size());
322  }
323  }
324  }
325 
326  bool empty() const noexcept { return MDeviceImages.empty(); }
327 
328  backend get_backend() const noexcept {
329  return MContext.get_platform().get_backend();
330  }
331 
332  context get_context() const noexcept { return MContext; }
333 
334  const std::vector<device> &get_devices() const noexcept { return MDevices; }
335 
336  std::vector<kernel_id> get_kernel_ids() const {
337  // Collect kernel ids from all device images, then remove duplicates
338 
339  std::vector<kernel_id> Result;
340  for (const device_image_plain &DeviceImage : MDeviceImages) {
341  const std::vector<kernel_id> &KernelIDs =
342  getSyclObjImpl(DeviceImage)->get_kernel_ids();
343 
344  Result.insert(Result.end(), KernelIDs.begin(), KernelIDs.end());
345  }
346  std::sort(Result.begin(), Result.end(), LessByNameComp{});
347 
348  auto NewIt = std::unique(Result.begin(), Result.end(), EqualByNameComp{});
349  Result.erase(NewIt, Result.end());
350 
351  return Result;
352  }
353 
354  kernel
355  get_kernel(const kernel_id &KernelID,
356  const std::shared_ptr<detail::kernel_bundle_impl> &Self) const {
357 
358  auto It = std::find_if(MDeviceImages.begin(), MDeviceImages.end(),
359  [&KernelID](const device_image_plain &DeviceImage) {
360  return DeviceImage.has_kernel(KernelID);
361  });
362 
363  if (MDeviceImages.end() == It)
364  throw sycl::exception(make_error_code(errc::invalid),
365  "The kernel bundle does not contain the kernel "
366  "identified by kernelId.");
367 
368  const std::shared_ptr<detail::device_image_impl> &DeviceImageImpl =
370 
371  RT::PiKernel Kernel = nullptr;
372  std::tie(Kernel, std::ignore) =
373  detail::ProgramManager::getInstance().getOrCreateKernel(
374  MContext, KernelID.get_name(), /*PropList=*/{},
375  DeviceImageImpl->get_program_ref());
376 
377  std::shared_ptr<kernel_impl> KernelImpl = std::make_shared<kernel_impl>(
378  Kernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self);
379 
380  return detail::createSyclObjFromImpl<kernel>(KernelImpl);
381  }
382 
383  bool has_kernel(const kernel_id &KernelID) const noexcept {
384  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
385  [&KernelID](const device_image_plain &DeviceImage) {
386  return DeviceImage.has_kernel(KernelID);
387  });
388  }
389 
390  bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept {
391  return std::any_of(
392  MDeviceImages.begin(), MDeviceImages.end(),
393  [&KernelID, &Dev](const device_image_plain &DeviceImage) {
394  return DeviceImage.has_kernel(KernelID, Dev);
395  });
396  }
397 
398  bool contains_specialization_constants() const noexcept {
399  return std::any_of(
400  MDeviceImages.begin(), MDeviceImages.end(),
401  [](const device_image_plain &DeviceImage) {
402  return getSyclObjImpl(DeviceImage)->has_specialization_constants();
403  });
404  }
405 
406  bool native_specialization_constant() const noexcept {
407  return std::all_of(MDeviceImages.begin(), MDeviceImages.end(),
408  [](const device_image_plain &DeviceImage) {
409  return getSyclObjImpl(DeviceImage)
410  ->all_specialization_constant_native();
411  });
412  }
413 
414  bool has_specialization_constant(const char *SpecName) const noexcept {
415  return std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
416  [SpecName](const device_image_plain &DeviceImage) {
417  return getSyclObjImpl(DeviceImage)
418  ->has_specialization_constant(SpecName);
419  });
420  }
421 
422  void set_specialization_constant_raw_value(const char *SpecName,
423  const void *Value,
424  size_t Size) noexcept {
425  if (has_specialization_constant(SpecName))
426  for (const device_image_plain &DeviceImage : MDeviceImages)
427  getSyclObjImpl(DeviceImage)
428  ->set_specialization_constant_raw_value(SpecName, Value);
429  else {
430  const auto *DataPtr = static_cast<const unsigned char *>(Value);
431  std::vector<unsigned char> &Val = MSpecConstValues[std::string{SpecName}];
432  Val.resize(Size);
433  Val.insert(Val.begin(), DataPtr, DataPtr + Size);
434  }
435  }
436 
437  void get_specialization_constant_raw_value(const char *SpecName,
438  void *ValueRet) const noexcept {
439  for (const device_image_plain &DeviceImage : MDeviceImages)
440  if (getSyclObjImpl(DeviceImage)->has_specialization_constant(SpecName)) {
441  getSyclObjImpl(DeviceImage)
442  ->get_specialization_constant_raw_value(SpecName, ValueRet);
443  return;
444  }
445 
446  // Specialization constant wasn't found in any of the device images,
447  // try to fetch value from kernel_bundle.
448  if (MSpecConstValues.count(std::string{SpecName}) != 0) {
449  const std::vector<unsigned char> &Val =
450  MSpecConstValues.at(std::string{SpecName});
451  auto *Dest = static_cast<unsigned char *>(ValueRet);
452  std::uninitialized_copy(Val.begin(), Val.end(), Dest);
453  return;
454  }
455 
456  assert(false &&
457  "get_specialization_constant_raw_value called for missing constant");
458  }
459 
460  bool is_specialization_constant_set(const char *SpecName) const noexcept {
461  bool SetInDevImg =
462  std::any_of(MDeviceImages.begin(), MDeviceImages.end(),
463  [SpecName](const device_image_plain &DeviceImage) {
464  return getSyclObjImpl(DeviceImage)
465  ->is_specialization_constant_set(SpecName);
466  });
467  return SetInDevImg || MSpecConstValues.count(std::string{SpecName}) != 0;
468  }
469 
470  const device_image_plain *begin() const { return MDeviceImages.data(); }
471 
472  const device_image_plain *end() const {
473  return MDeviceImages.data() + MDeviceImages.size();
474  }
475 
476  size_t size() const noexcept { return MDeviceImages.size(); }
477 
479  // All device images are expected to have the same state
480  return MDeviceImages.empty()
481  ? bundle_state::input
482  : detail::getSyclObjImpl(MDeviceImages[0])->get_state();
483  }
484 
485  const SpecConstMapT &get_spec_const_map_ref() const noexcept {
486  return MSpecConstValues;
487  }
488 
489  bool isInterop() const { return MIsInterop; }
490 
491 private:
492  context MContext;
493  std::vector<device> MDevices;
494  std::vector<device_image_plain> MDeviceImages;
495  // This map stores values for specialization constants, that are missing
496  // from any device image.
497  SpecConstMapT MSpecConstValues;
498  bool MIsInterop = false;
499 };
500 
501 } // namespace detail
502 } // namespace sycl
503 } // __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:476
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:81
PI_INVALID_OPERATION
@ PI_INVALID_OPERATION
Definition: pi.h:84
cl::sycl::detail::kernel_bundle_impl
Definition: kernel_bundle_impl.hpp:56
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:668
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:102
device.hpp
cl::sycl::detail::device_image_plain
Definition: kernel_bundle.hpp:70
cl::sycl::detail::kernel_bundle_impl::empty
bool empty() const noexcept
Definition: kernel_bundle_impl.hpp:326
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:485
cl::sycl::detail::kernel_bundle_impl::end
const device_image_plain * end() const
Definition: kernel_bundle_impl.hpp:472
context.hpp
cl::sycl::detail::LessByHash::operator()
bool operator()(const T &LHS, const T &RHS) const
Definition: kernel_bundle_impl.hpp:32
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:131
cl::sycl::detail::kernel_bundle_impl::has_specialization_constant
bool has_specialization_constant(const char *SpecName) const noexcept
Definition: kernel_bundle_impl.hpp:414
cl::sycl::detail::checkAllDevicesHaveAspect
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
Definition: kernel_bundle_impl.hpp:47
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:25
cl::sycl::detail::kernel_bundle_impl::isInterop
bool isInterop() const
Definition: kernel_bundle_impl.hpp:489
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:355
cl::sycl::detail::kernel_bundle_impl::get_devices
const std::vector< device > & get_devices() const noexcept
Definition: kernel_bundle_impl.hpp:334
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:578
cl::sycl::detail::kernel_bundle_impl::get_context
context get_context() const noexcept
Definition: kernel_bundle_impl.hpp:332
cl::sycl::detail::checkAllDevicesAreInContext
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
Definition: kernel_bundle_impl.hpp:37
cl::sycl::detail::kernel_bundle_impl::kernel_bundle_impl
kernel_bundle_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles)
Definition: kernel_bundle_impl.hpp:260
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:460
cl::sycl::detail::kernel_bundle_impl::begin
const device_image_plain * begin() const
Definition: kernel_bundle_impl.hpp:470
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:34
cl::sycl::detail::kernel_bundle_impl::has_kernel
bool has_kernel(const kernel_id &KernelID) const noexcept
Definition: kernel_bundle_impl.hpp:383
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:590
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:390
cl::sycl::kernel_id
Objects of the class identify kernel is some kernel_bundle related APIs.
Definition: kernel_bundle.hpp:38
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:437
cl::sycl::detail::tie
auto tie(Ts &... Args)
Definition: tuple.hpp:40
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:91
program_manager.hpp
cl::sycl::detail::kernel_bundle_impl::get_backend
backend get_backend() const noexcept
Definition: kernel_bundle_impl.hpp:328
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: kernel_bundle_impl.hpp:31
cl::sycl::detail::kernel_bundle_impl::get_bundle_state
bundle_state get_bundle_state() const
Definition: kernel_bundle_impl.hpp:478
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:447
cl::sycl::detail::kernel_bundle_impl::native_specialization_constant
bool native_specialization_constant() const noexcept
Definition: kernel_bundle_impl.hpp:406
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:182
backend_types.hpp
std
Definition: accessor.hpp:2358
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:422
cl::sycl::detail::kernel_bundle_impl::get_kernel_ids
std::vector< kernel_id > get_kernel_ids() const
Definition: kernel_bundle_impl.hpp:336
cl::sycl::context
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:35
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:398
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:106
__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:66