kernel.hpp

Go to the documentation of this file.

//==--------------- kernel.hpp --- SYCL kernel -----------------------------==//
//
// 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 <CL/sycl/detail/common.hpp>
#include <CL/sycl/detail/export.hpp>
#include <CL/sycl/detail/pi.h>
#include <CL/sycl/info/info_desc.hpp>
#include <CL/sycl/kernel_bundle_enums.hpp>
#include <CL/sycl/stl.hpp>
 
#include <memory>
 
__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
// Forward declaration
#ifdef __SYCL_INTERNAL_API
class program;
#endif
class context;
template <backend Backend> class backend_traits;
template <bundle_state State> class kernel_bundle;
 
namespace detail {
class kernel_impl;
 
class auto_name {};
 
template <typename Name, typename Type> struct get_kernel_name_t {
  using name = Name;
  static_assert(
      !std::is_same<Name, auto_name>::value,
      "No kernel name provided without -fsycl-unnamed-lambda enabled!");
};
 
#ifdef __SYCL_UNNAMED_LAMBDA__
template <typename Type> struct get_kernel_name_t<detail::auto_name, Type> {
  using name = Type;
};
#endif // __SYCL_UNNAMED_LAMBDA__
 
} // namespace detail
 
class __SYCL_EXPORT kernel {
public:
#ifdef __SYCL_INTERNAL_API
  kernel(cl_kernel ClKernel, const context &SyclContext);
#endif
 
  kernel(const kernel &RHS) = default;
 
  kernel(kernel &&RHS) = default;
 
  kernel &operator=(const kernel &RHS) = default;
 
  kernel &operator=(kernel &&RHS) = default;
 
  bool operator==(const kernel &RHS) const { return impl == RHS.impl; }
 
  bool operator!=(const kernel &RHS) const { return !operator==(RHS); }
 
#ifdef __SYCL_INTERNAL_API
  cl_kernel get() const;
#endif
 
  bool is_host() const;
 
  context get_context() const;
 
  backend get_backend() const noexcept;
 
  kernel_bundle<bundle_state::executable> get_kernel_bundle() const;
 
#ifdef __SYCL_INTERNAL_API
  program get_program() const;
#endif
 
  template <info::kernel param>
  typename info::param_traits<info::kernel, param>::return_type
  get_info() const;
 
  template <info::kernel_device_specific param>
  typename info::param_traits<info::kernel_device_specific, param>::return_type
  get_info(const device &Device) const;
 
  template <info::kernel_device_specific param>
  typename info::param_traits<info::kernel_device_specific, param>::return_type
  get_info(const device &Device,
           typename info::param_traits<info::kernel_device_specific,
                                       param>::input_type Value) const;
 
  template <info::kernel_work_group param>
  __SYCL2020_DEPRECATED("get_work_group_info() is deprecated, use SYCL 2020 "
                        "kernel_device_specific queries instead")
  typename info::param_traits<info::kernel_work_group, param>::return_type
      get_work_group_info(const device &Device) const;
 
  template <info::kernel_sub_group param>
  // clang-format off
  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
  typename info::param_traits<info::kernel_sub_group, param>::return_type
  get_sub_group_info(const device &Device) const;
  // clang-format on
 
  template <info::kernel_sub_group param>
  // clang-format off
  __SYCL_DEPRECATED("Use get_info with info::kernel_device_specific instead.")
  typename info::param_traits<info::kernel_sub_group, param>::return_type
  get_sub_group_info(const device &Device,
                     typename info::param_traits<info::kernel_sub_group,
                     param>::input_type Value) const;
  // clang-format on
 
  template <backend Backend>
  __SYCL_DEPRECATED("Use SYCL 2020 sycl::get_native free function")
  backend_return_t<Backend, kernel> get_native() const {
    return detail::pi::cast<backend_return_t<Backend, kernel>>(getNative());
  }
 
private:
  kernel(std::shared_ptr<detail::kernel_impl> Impl);
 
  pi_native_handle getNative() const;
 
  __SYCL_DEPRECATED("Use getNative() member function")
  pi_native_handle getNativeImpl() const;
 
  std::shared_ptr<detail::kernel_impl> impl;
 
  template <class Obj>
  friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
  template <class T>
  friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
};
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)
 
namespace std {
template <> struct hash<cl::sycl::kernel> {
  size_t operator()(const cl::sycl::kernel &Kernel) const {
    return hash<std::shared_ptr<cl::sycl::detail::kernel_impl>>()(
        cl::sycl::detail::getSyclObjImpl(Kernel));
  }
};
} // namespace std