DPC++ Runtime
Runtime libraries for oneAPI DPC++
kernel_impl.hpp
Go to the documentation of this file.
1 //==------- kernel_impl.hpp --- 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 
9 #pragma once
10 
12 #include <CL/sycl/detail/pi.h>
13 #include <CL/sycl/detail/pi.hpp>
14 #include <CL/sycl/device.hpp>
16 #include <CL/sycl/program.hpp>
17 #include <detail/context_impl.hpp>
18 #include <detail/device_impl.hpp>
19 #include <detail/kernel_info.hpp>
20 
21 #include <cassert>
22 #include <memory>
23 
25 namespace sycl {
26 namespace detail {
27 // Forward declaration
28 class program_impl;
29 class kernel_bundle_impl;
30 
31 using ContextImplPtr = std::shared_ptr<context_impl>;
32 using ProgramImplPtr = std::shared_ptr<program_impl>;
33 using KernelBundleImplPtr = std::shared_ptr<kernel_bundle_impl>;
34 class kernel_impl {
35 public:
45  kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context,
46  KernelBundleImplPtr KernelBundleImpl);
47 
61  kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
62  ProgramImplPtr ProgramImpl, bool IsCreatedFromSource,
63  KernelBundleImplPtr KernelBundleImpl);
64 
71  kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
72  DeviceImageImplPtr DeviceImageImpl,
73  KernelBundleImplPtr KernelBundleImpl);
74 
79  kernel_impl(ContextImplPtr Context, ProgramImplPtr ProgramImpl);
80 
81  // This section means the object is non-movable and non-copyable
82  // There is no need of move and copy constructors in kernel_impl.
83  // If they need to be added, piKernelRetain method for MKernel
84  // should be present.
85  kernel_impl(const kernel_impl &) = delete;
86  kernel_impl(kernel_impl &&) = delete;
87  kernel_impl &operator=(const kernel_impl &) = delete;
88  kernel_impl &operator=(kernel_impl &&) = delete;
89 
90  ~kernel_impl();
91 
99  cl_kernel get() const {
100  if (is_host()) {
101  throw invalid_object_error(
102  "This instance of kernel doesn't support OpenCL interoperability.",
104  }
106  return pi::cast<cl_kernel>(MKernel);
107  }
108 
112  bool is_host() const { return MContext->is_host(); }
113 
114  const plugin &getPlugin() const { return MContext->getPlugin(); }
115 
120  template <info::kernel param>
122  get_info() const;
123 
129  template <info::kernel_device_specific param>
131  get_info(const device &Device) const;
132 
139  template <info::kernel_device_specific param>
141  get_info(const device &Device,
143  param>::input_type Value) const;
144 
150  template <info::kernel_work_group param>
152  get_work_group_info(const device &Device) const;
153 
158  template <info::kernel_sub_group param>
160  get_sub_group_info(const device &Device) const;
161 
168  template <info::kernel_sub_group param>
170  get_sub_group_info(
171  const device &Device,
173  Value) const;
174 
178  RT::PiKernel &getHandleRef() { return MKernel; }
183  const RT::PiKernel &getHandleRef() const { return MKernel; }
184 
189  bool isCreatedFromSource() const;
190 
191  const DeviceImageImplPtr &getDeviceImage() const { return MDeviceImageImpl; }
192 
194  const plugin &Plugin = MContext->getPlugin();
195 
196  if (Plugin.getBackend() == backend::opencl)
197  Plugin.call<PiApiKind::piKernelRetain>(MKernel);
198 
199  pi_native_handle NativeKernel = 0;
200  Plugin.call<PiApiKind::piextKernelGetNativeHandle>(MKernel, &NativeKernel);
201 
202  return NativeKernel;
203  }
204 
205  KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; }
206 
207  bool isInterop() const { return MIsInterop; }
208 
209 private:
210  RT::PiKernel MKernel;
211  const ContextImplPtr MContext;
212  const ProgramImplPtr MProgramImpl;
213  bool MCreatedFromSource = true;
214  const DeviceImageImplPtr MDeviceImageImpl;
215  const KernelBundleImplPtr MKernelBundleImpl;
216  bool MIsInterop = false;
217 };
218 
219 template <info::kernel param>
221 kernel_impl::get_info() const {
222  if (is_host()) {
223  // TODO implement
224  assert(0 && "Not implemented");
225  }
226  return get_kernel_info<
228  param>::get(this->getHandleRef(), getPlugin());
229 }
230 
231 template <>
232 inline context kernel_impl::get_info<info::kernel::context>() const {
233  return createSyclObjFromImpl<context>(MContext);
234 }
235 
236 #ifdef __SYCL_INTERNAL_API
237 template <>
238 inline program kernel_impl::get_info<info::kernel::program>() const {
239  return createSyclObjFromImpl<program>(MProgramImpl);
240 }
241 #endif
242 
243 template <info::kernel_device_specific param>
244 inline typename info::param_traits<info::kernel_device_specific,
245  param>::return_type
246 kernel_impl::get_info(const device &Device) const {
247  if (is_host()) {
248  return get_kernel_device_specific_info_host<param>(Device);
249  }
252  param>::return_type,
253  param>::get(this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
254  getPlugin());
255 }
256 
257 template <info::kernel_device_specific param>
259  param>::return_type
260 kernel_impl::get_info(
261  const device &Device,
263  Value) const {
264  if (is_host()) {
265  throw runtime_error("Sub-group feature is not supported on HOST device.",
267  }
269  this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), Value,
270  getPlugin());
271 }
272 
273 template <info::kernel_work_group param>
275 kernel_impl::get_work_group_info(const device &Device) const {
276  return get_info<
278  Device);
279 }
280 
281 template <info::kernel_sub_group param>
283 kernel_impl::get_sub_group_info(const device &Device) const {
284  return get_info<
286  Device);
287 }
288 
289 template <info::kernel_sub_group param>
291 kernel_impl::get_sub_group_info(
292  const device &Device,
294  Value) const {
295  return get_info<
297  Device, Value);
298 }
299 
300 } // namespace detail
301 } // namespace sycl
302 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::kernel_impl::getNative
pi_native_handle getNative() const
Definition: kernel_impl.hpp:193
cl::sycl::detail::kernel_impl::get_kernel_bundle
KernelBundleImplPtr get_kernel_bundle() const
Definition: kernel_impl.hpp:205
cl::sycl::detail::kernel_impl::getDeviceImage
const DeviceImageImplPtr & getDeviceImage() const
Definition: kernel_impl.hpp:191
PI_INVALID_KERNEL
@ PI_INVALID_KERNEL
Definition: pi.h:89
cl::sycl::detail::pi::getPlugin
const plugin & getPlugin()
Definition: pi.cpp:511
pi.h
cl::sycl::detail::ContextImplPtr
std::shared_ptr< detail::context_impl > ContextImplPtr
Definition: memory_manager.hpp:32
context_impl.hpp
cl::sycl::detail::kernel_impl::isInterop
bool isInterop() const
Definition: kernel_impl.hpp:207
cl::sycl::info::param_traits
Definition: info_desc.hpp:310
piextKernelGetNativeHandle
pi_result piextKernelGetNativeHandle(pi_kernel kernel, pi_native_handle *nativeHandle)
Gets the native handle of a PI kernel object.
Definition: pi_esimd_emulator.cpp:1825
cl::sycl::detail::get_kernel_info
Definition: kernel_info.hpp:22
device.hpp
cl::sycl::detail::kernel_impl::getHandleRef
const RT::PiKernel & getHandleRef() const
Get a constant reference to a raw kernel object.
Definition: kernel_impl.hpp:183
cl::sycl::detail::KernelBundleImplPtr
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
Definition: kernel_bundle.hpp:132
sycl
Definition: invoke_simd.hpp:68
device_impl.hpp
cl::sycl::detail::kernel_impl
Definition: kernel_impl.hpp:34
pi.hpp
_pi_kernel
Implementation of a PI Kernel for CUDA.
Definition: pi_cuda.hpp:624
cl::sycl::detail::plugin::getBackend
backend getBackend(void) const
Definition: plugin.hpp:229
program.hpp
cl::sycl::detail::kernel_impl::get
cl_kernel get() const
Gets a valid OpenCL kernel handle.
Definition: kernel_impl.hpp:99
cl::sycl::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:35
cl::sycl::info::compatibility_param_traits
Definition: info_desc.hpp:312
cl::sycl::detail::plugin::call
void call(ArgsT... Args) const
Calls the API, traces the call, checks the result.
Definition: plugin.hpp:217
piKernelRetain
pi_result piKernelRetain(pi_kernel kernel)
Definition: pi_esimd_emulator.cpp:1381
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::plugin
The plugin class provides a unified interface to the underlying low-level runtimes for the device-agn...
Definition: plugin.hpp:90
pi_native_handle
uintptr_t pi_native_handle
Definition: pi.h:76
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(cl::sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
kernel_info.hpp
cl::sycl::detail::kernel_impl::getHandleRef
RT::PiKernel & getHandleRef()
Get a reference to a raw kernel object.
Definition: kernel_impl.hpp:178
cl::sycl::detail::get_kernel_device_specific_info
Definition: kernel_info.hpp:79
cl::sycl::detail::getSyclObjImpl
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
Definition: common.hpp:198
cl::sycl::detail::kernel_impl::is_host
bool is_host() const
Check if the associated SYCL context is a SYCL host context.
Definition: kernel_impl.hpp:112
cl::sycl::context
The context class represents a SYCL context on which kernel functions may be executed.
Definition: context.hpp:35
info_desc.hpp
common.hpp
cl::sycl::info::kernel_device_specific
kernel_device_specific
Definition: info_desc.hpp:265
cl::sycl::detail::ProgramImplPtr
std::shared_ptr< program_impl > ProgramImplPtr
Definition: kernel_impl.hpp:32
PI_INVALID_DEVICE
@ PI_INVALID_DEVICE
Definition: pi.h:94
__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
cl::sycl::detail::kernel_impl::getPlugin
const plugin & getPlugin() const
Definition: kernel_impl.hpp:114