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 {
18 namespace detail {
19 
20 kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context,
21  KernelBundleImplPtr KernelBundleImpl,
22  const KernelArgMask *ArgMask)
23  : kernel_impl(Kernel, Context,
24  std::make_shared<program_impl>(Context, Kernel),
25  /*IsCreatedFromSource*/ true, KernelBundleImpl, ArgMask) {
26  // Enable USM indirect access for interoperability kernels.
27  // Some PI Plugins (like OpenCL) require this call to enable USM
28  // For others, PI will turn this into a NOP.
30  MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE);
31 
32  // This constructor is only called in the interoperability kernel constructor.
33  MIsInterop = true;
34 }
35 
37  ProgramImplPtr ProgramImpl, bool IsCreatedFromSource,
38  KernelBundleImplPtr KernelBundleImpl,
39  const KernelArgMask *ArgMask)
40  : MKernel(Kernel), MContext(ContextImpl),
41  MProgramImpl(std::move(ProgramImpl)),
42  MCreatedFromSource(IsCreatedFromSource),
43  MKernelBundleImpl(std::move(KernelBundleImpl)),
44  MKernelArgMaskPtr{ArgMask} {
45 
46  RT::PiContext Context = nullptr;
47  // Using the plugin from the passed ContextImpl
49  MKernel, PI_KERNEL_INFO_CONTEXT, sizeof(Context), &Context, nullptr);
50  if (ContextImpl->getHandleRef() != Context)
51  throw sycl::invalid_parameter_error(
52  "Input context must be the same as the context of cl_kernel",
53  PI_ERROR_INVALID_CONTEXT);
54 
55  MIsInterop = MProgramImpl->isInterop();
56 }
57 
59  DeviceImageImplPtr DeviceImageImpl,
60  KernelBundleImplPtr KernelBundleImpl,
61  const KernelArgMask *ArgMask)
62  : MKernel(Kernel), MContext(std::move(ContextImpl)), MProgramImpl(nullptr),
63  MCreatedFromSource(false), MDeviceImageImpl(std::move(DeviceImageImpl)),
64  MKernelBundleImpl(std::move(KernelBundleImpl)),
65  MKernelArgMaskPtr{ArgMask} {
66 
67  // kernel_impl shared ownership of kernel handle
68  if (!is_host()) {
69  getPlugin()->call<PiApiKind::piKernelRetain>(MKernel);
70  }
71 
72  MIsInterop = MKernelBundleImpl->isInterop();
73 }
74 
76  : MContext(Context), MProgramImpl(std::move(ProgramImpl)) {}
77 
79  // TODO catch an exception and put it to list of asynchronous exceptions
80  if (!is_host()) {
81  getPlugin()->call<PiApiKind::piKernelRelease>(MKernel);
82  }
83 }
84 
86  // TODO it is not clear how to understand whether the SYCL kernel is created
87  // from source code or not when the SYCL kernel is created using
88  // the interoperability constructor.
89  // Here a strange case which does not work now:
90  // context Context;
91  // program Program(Context);
92  // Program.build_with_kernel_type<class A>();
93  // kernel FirstKernel= Program.get_kernel<class A>();
94  // cl_kernel ClKernel = FirstKernel.get();
95  // kernel SecondKernel = kernel(ClKernel, Context);
96  // clReleaseKernel(ClKernel);
97  // FirstKernel.isCreatedFromSource() != FirstKernel.isCreatedFromSource();
98  return MCreatedFromSource;
99 }
100 
101 bool kernel_impl::isBuiltInKernel(const device &Device) const {
102  auto BuiltInKernels = Device.get_info<info::device::built_in_kernel_ids>();
103  if (BuiltInKernels.empty())
104  return false;
105  std::string KernelName = get_info<info::kernel::function_name>();
106  return (std::any_of(
107  BuiltInKernels.begin(), BuiltInKernels.end(),
108  [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
109 }
110 
111 void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
112  if (MKernelBundleImpl->isInterop())
113  return;
114  auto Devices = MKernelBundleImpl->get_devices();
115  if (std::any_of(Devices.begin(), Devices.end(),
116  [this](device &Device) { return isBuiltInKernel(Device); }))
117  return;
118 
119  throw sycl::exception(
121  "info::kernel::num_args descriptor may only be used to query a kernel "
122  "that resides in a kernel bundle constructed using a backend specific"
123  "interoperability function or to query a device built-in kernel");
124 }
125 
126 } // namespace detail
127 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
128 } // namespace sycl
pi_bool
pi_uint32 pi_bool
Definition: pi.h:144
context_impl.hpp
sycl::_V1::detail::kernel_impl::isCreatedFromSource
bool isCreatedFromSource() const
Check if kernel was created from a program that had been created from source.
Definition: kernel_impl.cpp:85
sycl::_V1::detail::ContextImplPtr
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
Definition: event_impl.hpp:30
sycl::_V1::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:94
sycl::_V1::detail::DeviceImageImplPtr
std::shared_ptr< device_image_impl > DeviceImageImplPtr
Definition: kernel_bundle.hpp:72
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::detail::ProgramImplPtr
std::shared_ptr< program_impl > ProgramImplPtr
Definition: kernel_impl.hpp:32
sycl::_V1::detail::kernel_impl
Definition: kernel_impl.hpp:34
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
PI_USM_INDIRECT_ACCESS
@ PI_USM_INDIRECT_ACCESS
indicates that the kernel might access data through USM ptrs
Definition: pi.h:1428
sycl::_V1::detail::KernelBundleImplPtr
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
Definition: kernel_bundle.hpp:138
piKernelRelease
pi_result piKernelRelease(pi_kernel kernel)
Definition: pi_esimd_emulator.cpp:1438
sycl::_V1::detail::getPlugin
static const PluginPtr & getPlugin(backend Backend)
Definition: backend.cpp:32
program_impl.hpp
PI_KERNEL_INFO_CONTEXT
@ PI_KERNEL_INFO_CONTEXT
Definition: pi.h:406
piKernelSetExecInfo
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_esimd_emulator.cpp:2059
sycl::_V1::detail::KernelArgMask
std::vector< bool > KernelArgMask
Definition: kernel_arg_mask.hpp:16
sycl::_V1::detail::kernel_impl::~kernel_impl
~kernel_impl()
Definition: kernel_impl.cpp:78
sycl::_V1::detail::pi::PiContext
::pi_context PiContext
Definition: pi.hpp:130
piKernelRetain
pi_result piKernelRetain(pi_kernel kernel)
Definition: pi_esimd_emulator.cpp:1436
sycl::_V1::kernel_id
Objects of the class identify kernel is some kernel_bundle related APIs.
Definition: kernel_bundle.hpp:44
sycl::_V1::detail::kernel_impl::getPlugin
const PluginPtr & getPlugin() const
Definition: kernel_impl.hpp:117
sycl::_V1::device
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:49
sycl::_V1::errc::invalid
@ invalid
sycl::_V1::detail::pi::PiKernel
::pi_kernel PiKernel
Definition: pi.hpp:133
piKernelGetInfo
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_esimd_emulator.cpp:1421
std
Definition: accessor.hpp:3914
sycl::_V1::detail::program_impl
Definition: program_impl.hpp:38
kernel_bundle_impl.hpp
any_of
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
kernel_impl.hpp
sycl::_V1::detail::kernel_impl::is_host
bool is_host() const
Check if the associated SYCL context is a SYCL host context.
Definition: kernel_impl.hpp:115
PI_TRUE
const pi_bool PI_TRUE
Definition: pi.h:545
sycl::_V1::detail::kernel_impl::kernel_impl
kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context, KernelBundleImplPtr KernelBundleImpl, const KernelArgMask *ArgMask=nullptr)
Constructs a SYCL kernel instance from a PiKernel.
Definition: kernel_impl.cpp:20