DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_impl.cpp
Go to the documentation of this file.
1 //==------- kernel_impl.cpp --- SYCL kernel implementation -----------------==//
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 
11 #include <detail/kernel_impl.hpp>
12 #include <detail/program_impl.hpp>
13 
14 #include <memory>
15 
16 namespace sycl {
17 inline namespace _V1 {
18 namespace detail {
19 
21  ContextImplPtr Context,
22  KernelBundleImplPtr KernelBundleImpl,
23  const KernelArgMask *ArgMask)
24  : kernel_impl(Kernel, Context,
25  std::make_shared<program_impl>(Context, Kernel),
26  /*IsCreatedFromSource*/ true, KernelBundleImpl, ArgMask) {
27  // Enable USM indirect access for interoperability kernels.
28  // Some PI Plugins (like OpenCL) require this call to enable USM
29  // For others, PI will turn this into a NOP.
30  if (Context->getPlatformImpl()->supports_usm())
32  MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
33 
34  // This constructor is only called in the interoperability kernel constructor.
35  MIsInterop = true;
36 }
37 
39  ContextImplPtr ContextImpl, ProgramImplPtr ProgramImpl,
40  bool IsCreatedFromSource,
41  KernelBundleImplPtr KernelBundleImpl,
42  const KernelArgMask *ArgMask)
43  : MKernel(Kernel), MContext(ContextImpl),
44  MProgram(ProgramImpl->getHandleRef()),
45  MCreatedFromSource(IsCreatedFromSource),
46  MKernelBundleImpl(std::move(KernelBundleImpl)),
47  MKernelArgMaskPtr{ArgMask} {
48 
49  sycl::detail::pi::PiContext Context = nullptr;
50  // Using the plugin from the passed ContextImpl
52  MKernel, PI_KERNEL_INFO_CONTEXT, sizeof(Context), &Context, nullptr);
53  if (ContextImpl->getHandleRef() != Context)
54  throw sycl::invalid_parameter_error(
55  "Input context must be the same as the context of cl_kernel",
56  PI_ERROR_INVALID_CONTEXT);
57 
58  MIsInterop = ProgramImpl->isInterop();
59 }
60 
62  ContextImplPtr ContextImpl,
63  DeviceImageImplPtr DeviceImageImpl,
64  KernelBundleImplPtr KernelBundleImpl,
65  const KernelArgMask *ArgMask, PiProgram ProgramPI,
66  std::mutex *CacheMutex)
67  : MKernel(Kernel), MContext(std::move(ContextImpl)), MProgram(ProgramPI),
68  MCreatedFromSource(false), MDeviceImageImpl(std::move(DeviceImageImpl)),
69  MKernelBundleImpl(std::move(KernelBundleImpl)),
70  MKernelArgMaskPtr{ArgMask}, MCacheMutex{CacheMutex} {
71  MIsInterop = MKernelBundleImpl->isInterop();
72 }
73 
75  : MContext(Context), MProgram(ProgramImpl->getHandleRef()) {}
76 
78  // TODO catch an exception and put it to list of asynchronous exceptions
79  if (!is_host()) {
80  getPlugin()->call<PiApiKind::piKernelRelease>(MKernel);
81  }
82 }
83 
85  // TODO it is not clear how to understand whether the SYCL kernel is created
86  // from source code or not when the SYCL kernel is created using
87  // the interoperability constructor.
88  // Here a strange case which does not work now:
89  // context Context;
90  // program Program(Context);
91  // Program.build_with_kernel_type<class A>();
92  // kernel FirstKernel= Program.get_kernel<class A>();
93  // cl_kernel ClKernel = FirstKernel.get();
94  // kernel SecondKernel = kernel(ClKernel, Context);
95  // clReleaseKernel(ClKernel);
96  // FirstKernel.isCreatedFromSource() != FirstKernel.isCreatedFromSource();
97  return MCreatedFromSource;
98 }
99 
100 bool kernel_impl::isBuiltInKernel(const device &Device) const {
101  auto BuiltInKernels = Device.get_info<info::device::built_in_kernel_ids>();
102  if (BuiltInKernels.empty())
103  return false;
104  std::string KernelName = get_info<info::kernel::function_name>();
105  return (std::any_of(
106  BuiltInKernels.begin(), BuiltInKernels.end(),
107  [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
108 }
109 
110 void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
111  if (MKernelBundleImpl->isInterop())
112  return;
113  auto Devices = MKernelBundleImpl->get_devices();
114  if (std::any_of(Devices.begin(), Devices.end(),
115  [this](device &Device) { return isBuiltInKernel(Device); }))
116  return;
117 
118  throw sycl::exception(
120  "info::kernel::num_args descriptor may only be used to query a kernel "
121  "that resides in a kernel bundle constructed using a backend specific"
122  "interoperability function or to query a device built-in kernel");
123 }
124 
125 template <>
126 typename info::platform::version::return_type
127 kernel_impl::get_backend_info<info::platform::version>() const {
128  if (MContext->getBackend() != backend::opencl) {
130  "the info::platform::version info descriptor can "
131  "only be queried with an OpenCL backend");
132  }
133  auto Devices = MKernelBundleImpl->get_devices();
134  return Devices[0].get_platform().get_info<info::platform::version>();
135 }
136 
137 device select_device(DSelectorInvocableType DeviceSelectorInvocable,
138  std::vector<device> &Devices);
139 
140 template <>
141 typename info::device::version::return_type
142 kernel_impl::get_backend_info<info::device::version>() const {
143  if (MContext->getBackend() != backend::opencl) {
145  "the info::device::version info descriptor can only "
146  "be queried with an OpenCL backend");
147  }
148  auto Devices = MKernelBundleImpl->get_devices();
149  if (Devices.empty()) {
150  return "No available device";
151  }
152  // Use default selector to pick a device.
153  return select_device(default_selector_v, Devices)
154  .get_info<info::device::version>();
155 }
156 
157 template <>
158 typename info::device::backend_version::return_type
159 kernel_impl::get_backend_info<info::device::backend_version>() const {
160  if (MContext->getBackend() != backend::ext_oneapi_level_zero) {
162  "the info::device::backend_version info descriptor "
163  "can only be queried with a Level Zero backend");
164  }
165  return "";
166  // Currently The Level Zero backend does not define the value of this
167  // information descriptor and implementations are encouraged to return the
168  // empty string as per specification.
169 }
170 
171 } // namespace detail
172 } // namespace _V1
173 } // namespace sycl
kernel_impl(sycl::detail::pi::PiKernel Kernel, ContextImplPtr Context, KernelBundleImplPtr KernelBundleImpl, const KernelArgMask *ArgMask=nullptr)
Constructs a SYCL kernel instance from a PiKernel.
Definition: kernel_impl.cpp:20
bool is_host() const
Check if the associated SYCL context is a SYCL host context.
bool isCreatedFromSource() const
Check if kernel was created from a program that had been created from source.
Definition: kernel_impl.cpp:84
const PluginPtr & getPlugin() const
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:64
detail::is_device_info_desc< Param >::return_type get_info() const
Queries this SYCL device for information requested by the template parameter param.
Definition: device.hpp:223
Objects of the class identify kernel is some kernel_bundle related APIs.
class __SYCL2020_DEPRECATED("Host device is no longer supported.") host_selector int default_selector_v(const device &dev)
Selects SYCL host device.
::pi_program PiProgram
Definition: pi.hpp:137
std::shared_ptr< device_image_impl > DeviceImageImplPtr
std::vector< bool > KernelArgMask
std::shared_ptr< program_impl > ProgramImplPtr
Definition: kernel_impl.hpp:33
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
Definition: event_impl.hpp:32
std::function< int(const sycl::device &)> DSelectorInvocableType
device select_device(const DSelectorInvocableType &DeviceSelectorInvocable)
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:87
Definition: access.hpp:18
pi_result piKernelRelease(pi_kernel kernel)
Definition: pi_cuda.cpp:531
pi_uint32 pi_bool
Definition: pi.h:215
const pi_bool PI_TRUE
Definition: pi.h:684
@ PI_USM_INDIRECT_ACCESS
indicates that the kernel might access data through USM ptrs
Definition: pi.h:1603
pi_result piKernelSetExecInfo(pi_kernel kernel, pi_kernel_exec_info value_name, size_t param_value_size, const void *param_value)
API to set attributes controlling kernel execution.
Definition: pi_cuda.cpp:1053
@ PI_KERNEL_INFO_CONTEXT
Definition: pi.h:512
pi_result piKernelGetInfo(pi_kernel kernel, pi_kernel_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
Definition: pi_cuda.cpp:366
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept