DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_bundle.cpp
Go to the documentation of this file.
1 //==------- kernel_bundle.cpp - SYCL kernel_bundle and free functions ------==//
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 
14 
15 #include <cstddef>
16 #include <set>
17 #include <vector>
18 
19 namespace sycl {
20 inline namespace _V1 {
21 
22 kernel_id::kernel_id(const char *Name)
23  : impl(std::make_shared<detail::kernel_id_impl>(Name)) {}
24 
25 const char *kernel_id::get_name() const noexcept { return impl->get_name(); }
26 
27 namespace detail {
28 
32 
33 bool device_image_plain::has_kernel(const kernel_id &KernelID) const noexcept {
34  return impl->has_kernel(KernelID);
35 }
36 
37 bool device_image_plain::has_kernel(const kernel_id &KernelID,
38  const device &Dev) const noexcept {
39  return impl->has_kernel(KernelID, Dev);
40 }
41 
42 pi_native_handle device_image_plain::getNative() const {
43  return impl->getNative();
44 }
45 
49 
50 bool kernel_bundle_plain::empty() const noexcept { return impl->empty(); }
51 
52 backend kernel_bundle_plain::get_backend() const noexcept {
53  return impl->get_backend();
54 }
55 
56 context kernel_bundle_plain::get_context() const noexcept {
57  return impl->get_context();
58 }
59 
60 std::vector<device> kernel_bundle_plain::get_devices() const noexcept {
61  return impl->get_devices();
62 }
63 
64 std::vector<kernel_id> kernel_bundle_plain::get_kernel_ids() const {
65  return impl->get_kernel_ids();
66 }
67 
68 bool kernel_bundle_plain::contains_specialization_constants() const noexcept {
69  return impl->contains_specialization_constants();
70 }
71 
72 bool kernel_bundle_plain::native_specialization_constant() const noexcept {
73  return impl->native_specialization_constant();
74 }
75 
76 kernel kernel_bundle_plain::get_kernel(const kernel_id &KernelID) const {
77  return impl->get_kernel(KernelID, impl);
78 }
79 
80 const device_image_plain *kernel_bundle_plain::begin() const {
81  return impl->begin();
82 }
83 
84 const device_image_plain *kernel_bundle_plain::end() const {
85  return impl->end();
86 }
87 
88 bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID) const noexcept {
89  return impl->has_kernel(KernelID);
90 }
91 
92 bool kernel_bundle_plain::has_kernel(const kernel_id &KernelID,
93  const device &Dev) const noexcept {
94  return impl->has_kernel(KernelID, Dev);
95 }
96 
97 bool kernel_bundle_plain::has_specialization_constant_impl(
98  const char *SpecName) const noexcept {
99  return impl->has_specialization_constant(SpecName);
100 }
101 
102 void kernel_bundle_plain::set_specialization_constant_impl(
103  const char *SpecName, void *Value, size_t Size) noexcept {
104  impl->set_specialization_constant_raw_value(SpecName, Value, Size);
105 }
106 
107 void kernel_bundle_plain::get_specialization_constant_impl(
108  const char *SpecName, void *Value) const noexcept {
109  impl->get_specialization_constant_raw_value(SpecName, Value);
110 }
111 
112 bool kernel_bundle_plain::is_specialization_constant_set(
113  const char *SpecName) const noexcept {
114  return impl->is_specialization_constant_set(SpecName);
115 }
116 
117 bool kernel_bundle_plain::ext_oneapi_has_kernel(const std::string &name) {
118  return impl->ext_oneapi_has_kernel(name);
119 }
120 
121 kernel kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) {
122  return impl->ext_oneapi_get_kernel(name, impl);
123 }
124 
128 
129 const std::vector<device>
130 removeDuplicateDevices(const std::vector<device> &Devs) {
131  std::vector<device> UniqueDevices;
132 
133  // Building a new vector with unique elements and keep original order
134  std::unordered_set<device> UniqueDeviceSet;
135  for (const device &Dev : Devs)
136  if (UniqueDeviceSet.insert(Dev).second)
137  UniqueDevices.push_back(Dev);
138 
139  return UniqueDevices;
140 }
141 
143  return detail::ProgramManager::getInstance().getSYCLKernelID(
144  KernelName.data());
145 }
146 
148 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
149  bundle_state State) {
150  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, State);
151 }
152 
154 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
155  const std::vector<kernel_id> &KernelIDs,
156  bundle_state State) {
157  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, KernelIDs,
158  State);
159 }
160 
162 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
163  bundle_state State, const DevImgSelectorImpl &Selector) {
164  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, Selector,
165  State);
166 }
167 
170  const std::vector<device> &Devs) {
171  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
172 }
173 
174 std::shared_ptr<detail::kernel_bundle_impl>
175 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
176  bundle_state State) {
177  return std::make_shared<detail::kernel_bundle_impl>(Bundles, State);
178 }
179 
180 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
181  bundle_state State) {
182  // Check that all requested devices are associated with the context
183  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
184  if (Devs.empty() || !AllDevicesInTheContext)
185  throw sycl::exception(make_error_code(errc::invalid),
186  "Not all devices are associated with the context or "
187  "vector of devices is empty");
188 
189  if (bundle_state::input == State &&
190  !checkAllDevicesHaveAspect(Devs, aspect::online_compiler))
191  return false;
192  if (bundle_state::object == State &&
193  !checkAllDevicesHaveAspect(Devs, aspect::online_linker))
194  return false;
195 
196  const std::vector<device_image_plain> DeviceImages =
197  detail::ProgramManager::getInstance()
198  .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
199 
200  return (bool)DeviceImages.size();
201 }
202 
203 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
204  const std::vector<kernel_id> &KernelIds,
205  bundle_state State) {
206  // Check that all requested devices are associated with the context
207  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
208 
209  if (Devs.empty() || !AllDevicesInTheContext)
210  throw sycl::exception(make_error_code(errc::invalid),
211  "Not all devices are associated with the context or "
212  "vector of devices is empty");
213 
214  bool DeviceHasRequireAspectForState = true;
215  if (bundle_state::input == State) {
216  DeviceHasRequireAspectForState =
217  std::all_of(Devs.begin(), Devs.end(), [](const device &Dev) {
218  return Dev.has(aspect::online_compiler);
219  });
220  } else if (bundle_state::object == State) {
221  DeviceHasRequireAspectForState =
222  std::all_of(Devs.begin(), Devs.end(), [](const device &Dev) {
223  return Dev.has(aspect::online_linker);
224  });
225  }
226 
227  if (!DeviceHasRequireAspectForState)
228  return false;
229 
230  const std::vector<device_image_plain> DeviceImages =
231  detail::ProgramManager::getInstance()
232  .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
233 
234  std::set<kernel_id, LessByNameComp> CombinedKernelIDs;
235  for (const device_image_plain &DeviceImage : DeviceImages) {
236  const std::shared_ptr<device_image_impl> &DeviceImageImpl =
237  getSyclObjImpl(DeviceImage);
238 
239  CombinedKernelIDs.insert(DeviceImageImpl->get_kernel_ids_ptr()->begin(),
240  DeviceImageImpl->get_kernel_ids_ptr()->end());
241  }
242 
243  const bool AllKernelIDsRepresented =
244  std::all_of(KernelIds.begin(), KernelIds.end(),
245  [&CombinedKernelIDs](const kernel_id &KernelID) {
246  return CombinedKernelIDs.count(KernelID);
247  });
248 
249  return AllKernelIDsRepresented;
250 }
251 
252 std::shared_ptr<detail::kernel_bundle_impl>
254  const std::vector<device> &Devs, const property_list &PropList) {
255  return std::make_shared<detail::kernel_bundle_impl>(
256  InputBundle, Devs, PropList, bundle_state::object);
257 }
258 
259 std::shared_ptr<detail::kernel_bundle_impl>
260 link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
261  const std::vector<device> &Devs, const property_list &PropList) {
262  return std::make_shared<detail::kernel_bundle_impl>(ObjectBundles, Devs,
263  PropList);
264 }
265 
266 std::shared_ptr<detail::kernel_bundle_impl>
268  const std::vector<device> &Devs, const property_list &PropList) {
269  return std::make_shared<detail::kernel_bundle_impl>(
270  InputBundle, Devs, PropList, bundle_state::executable);
271 }
272 
273 // This function finds intersection of associated devices in common for all
274 // bundles
275 std::vector<sycl::device> find_device_intersection(
276  const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles) {
277  std::vector<sycl::device> IntersectDevices;
278  std::vector<unsigned int> DevsCounters;
279  std::map<device, unsigned int, LessByHash<device>> DevCounters;
280  for (const sycl::kernel_bundle<bundle_state::object> &ObjectBundle :
281  ObjectBundles)
282  // Increment counter in "DevCounters" each time a device is seen
283  for (const sycl::device &Device : ObjectBundle.get_devices())
284  DevCounters[Device]++;
285 
286  // If some device counter is less than ObjectBundles.size() then some bundle
287  // doesn't have it - do not add such a device to the final result
288  for (const std::pair<const device, unsigned int> &It : DevCounters)
289  if (ObjectBundles.size() == It.second)
290  IntersectDevices.push_back(It.first);
291 
292  return IntersectDevices;
293 }
294 
295 } // namespace detail
296 
300 
301 std::vector<kernel_id> get_kernel_ids() {
302  return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
303 }
304 
305 bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
306  if (KernelIDs.empty())
307  return true;
308  // TODO: also need to check that the architecture specified by the
309  // "-fsycl-targets" flag matches the device when we are able to get the
310  // device's arch.
311  auto doesImageTargetMatchDevice = [](const device &Dev,
312  const detail::RTDeviceBinaryImage &Img) {
313  const char *Target = Img.getRawData().DeviceTargetSpec;
314  auto BE = Dev.get_backend();
315  // ESIMD emulator backend is only compatible with esimd kernels.
316  if (BE == sycl::backend::ext_intel_esimd_emulator) {
317  pi_device_binary_property Prop = Img.getProperty("isEsimdImage");
318  return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0));
319  }
320  if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) {
321  return (BE == sycl::backend::opencl ||
322  BE == sycl::backend::ext_oneapi_level_zero);
323  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) ==
324  0) {
325  return Dev.is_cpu();
326  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN) ==
327  0) {
328  return Dev.is_gpu() && (BE == sycl::backend::opencl ||
329  BE == sycl::backend::ext_oneapi_level_zero);
330  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA) ==
331  0) {
332  return Dev.is_accelerator();
333  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64) == 0) {
334  return BE == sycl::backend::ext_oneapi_cuda;
335  } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN) == 0) {
336  return BE == sycl::backend::ext_oneapi_hip;
337  }
338 
339  return false;
340  };
341 
342  // One kernel may be contained in several binary images depending on the
343  // number of targets. This kernel is compatible with the device if there is
344  // at least one image (containing this kernel) whose aspects are supported by
345  // the device and whose target matches the device.
346  for (const auto &KernelID : KernelIDs) {
347  std::set<detail::RTDeviceBinaryImage *> BinImages =
348  detail::ProgramManager::getInstance().getRawDeviceImages({KernelID});
349 
350  if (std::none_of(BinImages.begin(), BinImages.end(),
351  [&](const detail::RTDeviceBinaryImage *Img) {
352  return doesDevSupportDeviceRequirements(Dev, *Img) &&
353  doesImageTargetMatchDevice(Dev, *Img);
354  }))
355  return false;
356  }
357 
358  return true;
359 }
360 
362 // * kernel_compiler extension *
364 namespace ext::oneapi::experimental {
365 
369 
371 // syclex::is_source_kernel_bundle_supported
374  // Support is limited to the opencl and level_zero backends.
375  bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
376  (BE == sycl::backend::opencl);
377  if (BE_Acceptable) {
378  // At the moment, OpenCL and SPIR-V are the only supported languages.
379  if (Language == source_language::opencl) {
381  } else if (Language == source_language::spirv) {
382  return true;
383  }
384  }
385 
386  // otherwise
387  return false;
388 }
389 
391 // syclex::create_kernel_bundle_from_source
393 
395  source_language Language,
396  const std::string &Source) {
397  // TODO: if we later support a "reason" why support isn't present
398  // (like a missing shared library etc.) it'd be nice to include it in
399  // the exception message here.
400  backend BE = SyclContext.get_backend();
401  if (!is_source_kernel_bundle_supported(BE, Language))
402  throw sycl::exception(make_error_code(errc::invalid),
403  "kernel_bundle creation from source not supported");
404 
405  std::shared_ptr<kernel_bundle_impl> KBImpl =
406  std::make_shared<kernel_bundle_impl>(SyclContext, Language, Source);
407  return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
408 }
409 
410 source_kb
412  source_language Language,
413  const std::vector<std::byte> &Bytes) {
414  backend BE = SyclContext.get_backend();
415  if (!is_source_kernel_bundle_supported(BE, Language))
416  throw sycl::exception(make_error_code(errc::invalid),
417  "kernel_bundle creation from source not supported");
418 
419  std::shared_ptr<kernel_bundle_impl> KBImpl =
420  std::make_shared<kernel_bundle_impl>(SyclContext, Language, Bytes);
421  return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
422 }
423 
425 // syclex::detail::build_from_source(source_kb) => exe_kb
427 namespace detail {
428 
430  const std::vector<device> &Devices,
431  const std::vector<std::string> &BuildOptions,
432  std::string *LogPtr) {
433  std::vector<device> UniqueDevices =
435  std::shared_ptr<kernel_bundle_impl> sourceImpl = getSyclObjImpl(SourceKB);
436  std::shared_ptr<kernel_bundle_impl> KBImpl =
437  sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr);
438  return sycl::detail::createSyclObjFromImpl<exe_kb>(KBImpl);
439 }
440 
441 } // namespace detail
442 } // namespace ext::oneapi::experimental
443 
444 } // namespace _V1
445 } // namespace sycl
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:50
backend get_backend() const noexcept
Returns the backend associated with this context.
Definition: context.cpp:146
const char * data() const noexcept
Definition: string_view.hpp:38
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:64
bool is_accelerator() const
Check if device is an accelerator device.
Definition: device.cpp:83
bool is_gpu() const
Check if device is a GPU device.
Definition: device.cpp:81
static std::vector< device > get_devices(info::device_type deviceType=info::device_type::all)
Query available SYCL devices.
Definition: device.cpp:51
backend get_backend() const noexcept
Returns the backend associated with this device.
Definition: device.cpp:215
bool is_cpu() const
Check if device is a CPU device.
Definition: device.cpp:79
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.
Provides an abstraction of a SYCL kernel.
Definition: kernel.hpp:76
Objects of the property_list class are containers for the SYCL properties.
std::vector< sycl::device > find_device_intersection(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles)
std::shared_ptr< detail::kernel_bundle_impl > build_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
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)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: impl_utils.hpp:30
const std::vector< device > removeDuplicateDevices(const std::vector< device > &Devs)
detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
std::shared_ptr< detail::kernel_bundle_impl > compile_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
std::shared_ptr< detail::kernel_bundle_impl > join_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
kernel_id get_kernel_id_impl(string_view KernelName)
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
bool has_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
detail::KernelBundleImplPtr get_empty_interop_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs)
std::shared_ptr< detail::kernel_bundle_impl > link_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, const std::vector< device > &Devs, const property_list &PropList)
exe_kb build_from_source(source_kb &SourceKB, const std::vector< device > &Devices, const std::vector< std::string > &BuildOptions, std::string *LogPtr)
bool is_source_kernel_bundle_supported(backend BE, source_language Language)
kernel_bundle< bundle_state::ext_oneapi_source > create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::string &Source)
sycl::detail::kernel_bundle_impl kernel_bundle_impl
kernel_bundle< sycl::bundle_state::ext_oneapi_source > source_kb
std::vector< kernel_id > get_kernel_ids()
bool is_compatible(const std::vector< kernel_id > &KernelIDs, const device &Dev)
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
uintptr_t pi_native_handle
Definition: pi.h:243
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64
Device-specific binary images produced from SPIR-V 64-bit <-> various "spir64_*" triples for specific...
Definition: pi.h:1012
#define __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64
PTX 64-bit image <-> "nvptx64", 64-bit NVIDIA PTX device.
Definition: pi.h:1016
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64
SPIR-V 64-bit image <-> "spir64", 64-bit OpenCL device.
Definition: pi.h:1009
#define __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN
Definition: pi.h:1017
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN
Definition: pi.h:1013
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA
Definition: pi.h:1014
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