DPC++ Runtime
Runtime libraries for oneAPI Data Parallel C++
kernel_desc.hpp
Go to the documentation of this file.
1 //==----------------------- kernel_desc.hpp --------------------------------==//
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 
11 // This header file must not include any standard C++ header files.
12 
15 
17 namespace sycl {
18 namespace detail {
19 
20 #ifndef __SYCL_DEVICE_ONLY__
21 #define _Bool bool
22 #endif
23 
24 // kernel parameter kinds
25 enum class kernel_param_kind_t {
26  kind_accessor = 0,
27  kind_std_layout = 1, // standard layout object parameters
28  kind_sampler = 2,
29  kind_pointer = 3,
31  kind_stream = 5,
32  kind_invalid = 0xf, // not a valid kernel kind
33 };
34 
35 // describes a kernel parameter
37  // parameter kind
39  // kind == kind_std_layout
40  // parameter size in bytes (includes padding for structs)
41  // kind == kind_accessor
42  // access target; possible access targets are defined in access/access.hpp
43  int info;
44  // offset of the captured value of the parameter in the lambda or function
45  // object
46  int offset;
47 };
48 
49 // Translates specialization constant type to its name.
50 template <class Name> struct SpecConstantInfo {
51  static constexpr const char *getName() { return ""; }
52 };
53 
54 // This guard is needed because the libsycl.so can be compiled with C++ <=14
55 // while the code requires C++17. This code is not supposed to be used by the
56 // libsycl.so so it should not be a problem.
57 #if __cplusplus >= 201703L
58 // Translates SYCL 2020 `specialization_id` to a unique symbolic identifier.
59 // There are no primary definition, only specializations in the integration
60 // footer.
61 template <auto &SpecName> const char *get_spec_constant_symbolic_ID_impl();
62 // Wrapper is needed to delay instantiation of
63 // 'get_spec_constant_symbolic_ID_impl' until after we have encountered all
64 // specializations for it generated by the compiler in the integration footer.
65 // Definition in spec_const_integration.hpp.
66 template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
67 #endif
68 
69 #ifndef __SYCL_UNNAMED_LAMBDA__
70 template <class KernelNameType> struct KernelInfo {
71  static constexpr unsigned getNumParams() { return 0; }
72  static const kernel_param_desc_t &getParamDesc(int) {
73  static kernel_param_desc_t Dummy;
74  return Dummy;
75  }
76  static constexpr const char *getName() { return ""; }
77  static constexpr bool isESIMD() { return 0; }
78 };
79 #else
80 template <char...> struct KernelInfoData {
81  static constexpr unsigned getNumParams() { return 0; }
82  static const kernel_param_desc_t &getParamDesc(int Idx) {
83  static kernel_param_desc_t Dummy;
84  return Dummy;
85  }
86  static constexpr const char *getName() { return ""; }
87  static constexpr bool isESIMD() { return 0; }
88 };
89 
90 // C++14 like index_sequence and make_index_sequence
91 // not needed C++14 members (value_type, size) not implemented
92 template <class T, T...> struct integer_sequence {};
93 template <unsigned long long... I>
94 using index_sequence = integer_sequence<unsigned long long, I...>;
95 template <unsigned long long N>
96 using make_index_sequence =
97  __make_integer_seq<integer_sequence, unsigned long long, N>;
98 
99 template <typename T> struct KernelInfoImpl {
100 private:
101  static constexpr auto n = __builtin_sycl_unique_stable_name(T);
102  template <unsigned long long... I>
103  static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
104  return {};
105  }
106 
107 public:
108  using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{}));
109 };
110 
111 // For named kernels, this structure is specialized in the integration header.
112 // For unnamed kernels, KernelInfoData is specialized in the integration header,
113 // and this picks it up via the KernelInfoImpl. For non-existent kernels, this
114 // will also pick up a KernelInfoData (as SubKernelInfo) via KernelInfoImpl, but
115 // it will instead get the unspecialized case, defined above.
116 template <class KernelNameType> struct KernelInfo {
117  using SubKernelInfo = typename KernelInfoImpl<KernelNameType>::type;
118  static constexpr unsigned getNumParams() {
119  return SubKernelInfo::getNumParams();
120  }
121  static const kernel_param_desc_t &getParamDesc(int Idx) {
122  return SubKernelInfo::getParamDesc(Idx);
123  }
124  static constexpr const char *getName() { return SubKernelInfo::getName(); }
125  static constexpr bool isESIMD() { return SubKernelInfo::isESIMD(); }
126 };
127 #endif //__SYCL_UNNAMED_LAMBDA__
128 
129 } // namespace detail
130 } // namespace sycl
131 } // __SYCL_INLINE_NAMESPACE(cl)
type
cl::sycl::detail::KernelInfo::getName
static constexpr const char * getName()
Definition: kernel_desc.hpp:76
T
defines_elementary.hpp
cl::sycl::ext::intel::experimental::type
type
Definition: fpga_utils.hpp:22
cl::sycl::detail::SpecConstantInfo::getName
static constexpr const char * getName()
Definition: kernel_desc.hpp:51
cl::sycl::detail::kernel_param_kind_t::kind_sampler
@ kind_sampler
detail
Definition: pi_opencl.cpp:86
cl::sycl::detail::kernel_param_kind_t::kind_pointer
@ kind_pointer
export.hpp
cl::sycl::detail::KernelInfo::isESIMD
static constexpr bool isESIMD()
Definition: kernel_desc.hpp:77
cl::sycl::detail::KernelInfo::getNumParams
static constexpr unsigned getNumParams()
Definition: kernel_desc.hpp:71
cl::sycl::detail::kernel_param_kind_t::kind_std_layout
@ kind_std_layout
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::detail::kernel_param_kind_t::kind_specialization_constants_buffer
@ kind_specialization_constants_buffer
cl::sycl::detail::SpecConstantInfo
Definition: kernel_desc.hpp:50
cl::sycl::detail::KernelInfo
Definition: kernel_desc.hpp:70
cl::sycl::detail::kernel_param_kind_t::kind_stream
@ kind_stream
cl::sycl::detail::kernel_param_desc_t::offset
int offset
Definition: kernel_desc.hpp:46
cl::sycl::detail::kernel_param_kind_t::kind_accessor
@ kind_accessor
cl::sycl::detail::kernel_param_kind_t::kind_invalid
@ kind_invalid
cl::sycl::detail::kernel_param_kind_t
kernel_param_kind_t
Definition: kernel_desc.hpp:25
cl::sycl::detail::kernel_param_desc_t
Definition: kernel_desc.hpp:36
cl::sycl::detail::kernel_param_desc_t::kind
kernel_param_kind_t kind
Definition: kernel_desc.hpp:38
cl::sycl::detail::kernel_param_desc_t::info
int info
Definition: kernel_desc.hpp:43
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12
cl::sycl::detail::KernelInfo::getParamDesc
static const kernel_param_desc_t & getParamDesc(int)
Definition: kernel_desc.hpp:72