kernel_impl.hpp

Go to the documentation of this file.

//==------- kernel_impl.hpp --- SYCL kernel implementation -----------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <detail/context_impl.hpp>
#include <detail/device_impl.hpp>
#include <detail/kernel_info.hpp>
#include <sycl/detail/common.hpp>
#include <sycl/detail/pi.h>
#include <sycl/detail/pi.hpp>
#include <sycl/device.hpp>
#include <sycl/info/info_desc.hpp>
 
#include <cassert>
#include <memory>
 
namespace sycl {
__SYCL_INLINE_VER_NAMESPACE(_V1) {
namespace detail {
// Forward declaration
class program_impl;
class kernel_bundle_impl;
 
using ContextImplPtr = std::shared_ptr<context_impl>;
using ProgramImplPtr = std::shared_ptr<program_impl>;
using KernelBundleImplPtr = std::shared_ptr<kernel_bundle_impl>;
class kernel_impl {
public:
  kernel_impl(RT::PiKernel Kernel, ContextImplPtr Context,
              KernelBundleImplPtr KernelBundleImpl);
 
  kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
              ProgramImplPtr ProgramImpl, bool IsCreatedFromSource,
              KernelBundleImplPtr KernelBundleImpl);
 
  kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
              DeviceImageImplPtr DeviceImageImpl,
              KernelBundleImplPtr KernelBundleImpl);
 
  kernel_impl(ContextImplPtr Context, ProgramImplPtr ProgramImpl);
 
  // This section means the object is non-movable and non-copyable
  // There is no need of move and copy constructors in kernel_impl.
  // If they need to be added, piKernelRetain method for MKernel
  // should be present.
  kernel_impl(const kernel_impl &) = delete;
  kernel_impl(kernel_impl &&) = delete;
  kernel_impl &operator=(const kernel_impl &) = delete;
  kernel_impl &operator=(kernel_impl &&) = delete;
 
  ~kernel_impl();
 
  cl_kernel get() const {
    if (is_host()) {
      throw invalid_object_error(
          "This instance of kernel doesn't support OpenCL interoperability.",
          PI_ERROR_INVALID_KERNEL);
    }
    getPlugin().call<PiApiKind::piKernelRetain>(MKernel);
    return pi::cast<cl_kernel>(MKernel);
  }
 
  bool is_host() const { return MContext->is_host(); }
 
  const plugin &getPlugin() const { return MContext->getPlugin(); }
 
  template <typename Param> typename Param::return_type get_info() const;
 
  template <typename Param>
  typename Param::return_type get_info(const device &Device) const;
 
  template <typename Param>
  typename Param::return_type get_info(const device &Device,
                                       const range<3> &WGSize) const;
 
  RT::PiKernel &getHandleRef() { return MKernel; }
  const RT::PiKernel &getHandleRef() const { return MKernel; }
 
  bool isCreatedFromSource() const;
 
  const DeviceImageImplPtr &getDeviceImage() const { return MDeviceImageImpl; }
 
  pi_native_handle getNative() const {
    const plugin &Plugin = MContext->getPlugin();
 
    if (Plugin.getBackend() == backend::opencl)
      Plugin.call<PiApiKind::piKernelRetain>(MKernel);
 
    pi_native_handle NativeKernel = 0;
    Plugin.call<PiApiKind::piextKernelGetNativeHandle>(MKernel, &NativeKernel);
 
    return NativeKernel;
  }
 
  KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; }
 
  bool isInterop() const { return MIsInterop; }
 
  ProgramImplPtr getProgramImpl() const { return MProgramImpl; }
 
  std::mutex &getNoncacheableEnqueueMutex() {
    return MNoncacheableEnqueueMutex;
  }
 
private:
  RT::PiKernel MKernel;
  const ContextImplPtr MContext;
  const ProgramImplPtr MProgramImpl;
  bool MCreatedFromSource = true;
  const DeviceImageImplPtr MDeviceImageImpl;
  const KernelBundleImplPtr MKernelBundleImpl;
  bool MIsInterop = false;
  std::mutex MNoncacheableEnqueueMutex;
};
 
template <typename Param>
inline typename Param::return_type kernel_impl::get_info() const {
  static_assert(is_kernel_info_desc<Param>::value,
                "Invalid kernel information descriptor");
  if (is_host()) {
    // TODO implement
    assert(0 && "Not implemented");
  }
  return get_kernel_info<Param>(this->getHandleRef(), getPlugin());
}
 
template <>
inline context kernel_impl::get_info<info::kernel::context>() const {
  return createSyclObjFromImpl<context>(MContext);
}
 
template <typename Param>
inline typename Param::return_type
kernel_impl::get_info(const device &Device) const {
  if (is_host()) {
    return get_kernel_device_specific_info_host<Param>(Device);
  }
  return get_kernel_device_specific_info<Param>(
      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(),
      getPlugin());
}
 
template <typename Param>
inline typename Param::return_type
kernel_impl::get_info(const device &Device,
                      const sycl::range<3> &WGSize) const {
  if (is_host()) {
    throw runtime_error("Sub-group feature is not supported on HOST device.",
                        PI_ERROR_INVALID_DEVICE);
  }
  return get_kernel_device_specific_info_with_input<Param>(
      this->getHandleRef(), getSyclObjImpl(Device)->getHandleRef(), WGSize,
      getPlugin());
}
 
} // namespace detail
} // __SYCL_INLINE_VER_NAMESPACE(_V1)
} // namespace sycl