DPC++ Runtime
Runtime libraries for oneAPI DPC++
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 
14 #include <sycl/detail/export.hpp>
15 
16 namespace sycl {
18 namespace detail {
19 
20 #ifndef __SYCL_DEVICE_ONLY__
21 #define _Bool bool
22 #endif
23 
24 // As stated above, this header file cannot include any of the standard C++
25 // headers. We need int64_t. Here we are matching the exact definition used by
26 // the SemaSYCL version of kernel_desc.hpp in the FE.
27 template <bool Cond, typename TrueT, typename FalseT> struct conditional {
28  using type = TrueT;
29 };
30 template <typename TrueT, typename FalseT>
31 struct conditional<false, TrueT, FalseT> {
32  using type = FalseT;
33 };
34 using int64_t = conditional<sizeof(long) == 8, long, long long>::type;
35 
36 // kernel parameter kinds
37 enum class kernel_param_kind_t {
38  kind_accessor = 0,
39  kind_std_layout = 1, // standard layout object parameters
40  kind_sampler = 2,
41  kind_pointer = 3,
43  kind_stream = 5,
44  kind_invalid = 0xf, // not a valid kernel kind
45 };
46 
47 // describes a kernel parameter
49  // parameter kind
51  // kind == kind_std_layout
52  // parameter size in bytes (includes padding for structs)
53  // kind == kind_accessor
54  // access target; possible access targets are defined in access/access.hpp
55  int info;
56  // offset of the captured value of the parameter in the lambda or function
57  // object
58  int offset;
59 };
60 
61 // Translates specialization constant type to its name.
62 template <class Name> struct SpecConstantInfo {
63  static constexpr const char *getName() { return ""; }
64 };
65 
66 // Translates SYCL 2020 `specialization_id` to a unique symbolic identifier.
67 // There are no primary definition, only specializations in the integration
68 // footer.
69 template <auto &SpecName> const char *get_spec_constant_symbolic_ID_impl();
70 // Wrapper is needed to delay instantiation of
71 // 'get_spec_constant_symbolic_ID_impl' until after we have encountered all
72 // specializations for it generated by the compiler in the integration footer.
73 // Definition in spec_const_integration.hpp.
74 template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
75 
76 #ifndef __SYCL_UNNAMED_LAMBDA__
77 template <class KernelNameType> struct KernelInfo {
78  static constexpr unsigned getNumParams() { return 0; }
79  static const kernel_param_desc_t &getParamDesc(int) {
80  static kernel_param_desc_t Dummy;
81  return Dummy;
82  }
83  static constexpr const char *getName() { return ""; }
84  static constexpr bool isESIMD() { return 0; }
85  static constexpr const char *getFileName() { return ""; }
86  static constexpr const char *getFunctionName() { return ""; }
87  static constexpr unsigned getLineNumber() { return 0; }
88  static constexpr unsigned getColumnNumber() { return 0; }
89  static constexpr int64_t getKernelSize() { return 0; }
90 };
91 #else
92 template <char...> struct KernelInfoData {
93  static constexpr unsigned getNumParams() { return 0; }
94  static const kernel_param_desc_t &getParamDesc(int Idx) {
95  static kernel_param_desc_t Dummy;
96  return Dummy;
97  }
98  static constexpr const char *getName() { return ""; }
99  static constexpr bool isESIMD() { return 0; }
100  static constexpr const char *getFileName() { return ""; }
101  static constexpr const char *getFunctionName() { return ""; }
102  static constexpr unsigned getLineNumber() { return 0; }
103  static constexpr unsigned getColumnNumber() { return 0; }
104  static constexpr int64_t getKernelSize() { return 0; }
105 };
106 
107 // C++14 like index_sequence and make_index_sequence
108 // not needed C++14 members (value_type, size) not implemented
109 template <class T, T...> struct integer_sequence {};
110 template <unsigned long long... I>
111 using index_sequence = integer_sequence<unsigned long long, I...>;
112 template <unsigned long long N>
113 using make_index_sequence =
114  __make_integer_seq<integer_sequence, unsigned long long, N>;
115 
116 template <typename T> struct KernelInfoImpl {
117 private:
118  static constexpr auto n = __builtin_sycl_unique_stable_name(T);
119  template <unsigned long long... I>
120  static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
121  return {};
122  }
123 
124 public:
125  using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{}));
126 };
127 
128 // For named kernels, this structure is specialized in the integration header.
129 // For unnamed kernels, KernelInfoData is specialized in the integration header,
130 // and this picks it up via the KernelInfoImpl. For non-existent kernels, this
131 // will also pick up a KernelInfoData (as SubKernelInfo) via KernelInfoImpl, but
132 // it will instead get the unspecialized case, defined above.
133 template <class KernelNameType> struct KernelInfo {
134  using SubKernelInfo = typename KernelInfoImpl<KernelNameType>::type;
135  static constexpr unsigned getNumParams() {
136  return SubKernelInfo::getNumParams();
137  }
138  static const kernel_param_desc_t &getParamDesc(int Idx) {
139  return SubKernelInfo::getParamDesc(Idx);
140  }
141  static constexpr const char *getName() { return SubKernelInfo::getName(); }
142  static constexpr bool isESIMD() { return SubKernelInfo::isESIMD(); }
143  static constexpr const char *getFileName() { return ""; }
144  static constexpr const char *getFunctionName() { return ""; }
145  static constexpr unsigned getLineNumber() { return 0; }
146  static constexpr unsigned getColumnNumber() { return 0; }
147  static constexpr int64_t getKernelSize() {
148  return SubKernelInfo::getKernelSize();
149  }
150 };
151 #endif //__SYCL_UNNAMED_LAMBDA__
152 
153 } // namespace detail
154 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
155 } // namespace sycl
sycl::_V1::detail::KernelInfo::isESIMD
static constexpr bool isESIMD()
Definition: kernel_desc.hpp:84
sycl::_V1::detail::kernel_param_kind_t::kind_stream
@ kind_stream
sycl::_V1::detail::kernel_param_kind_t::kind_std_layout
@ kind_std_layout
sycl::_V1::detail::kernel_param_desc_t
Definition: kernel_desc.hpp:48
sycl::_V1::detail::kernel_param_kind_t
kernel_param_kind_t
Definition: kernel_desc.hpp:37
sycl::_V1::detail::conditional
Definition: kernel_desc.hpp:27
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::detail::get_spec_constant_symbolic_ID
const char * get_spec_constant_symbolic_ID()
Definition: spec_const_integration.hpp:19
sycl::_V1::detail::int64_t
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:34
sycl::_V1::detail::kernel_param_kind_t::kind_accessor
@ kind_accessor
sycl::_V1::detail::KernelInfo::getLineNumber
static constexpr unsigned getLineNumber()
Definition: kernel_desc.hpp:87
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::detail::kernel_param_desc_t::info
int info
Definition: kernel_desc.hpp:55
export.hpp
sycl::_V1::detail::conditional< false, TrueT, FalseT >::type
FalseT type
Definition: kernel_desc.hpp:32
defines_elementary.hpp
sycl::_V1::detail::SpecConstantInfo
Definition: kernel_desc.hpp:62
sycl::_V1::detail::kernel_param_desc_t::kind
kernel_param_kind_t kind
Definition: kernel_desc.hpp:50
sycl::_V1::detail::KernelInfo::getFunctionName
static constexpr const char * getFunctionName()
Definition: kernel_desc.hpp:86
sycl::_V1::detail::conditional::type
TrueT type
Definition: kernel_desc.hpp:28
sycl::_V1::detail::KernelInfo::getName
static constexpr const char * getName()
Definition: kernel_desc.hpp:83
sycl::_V1::detail::kernel_param_kind_t::kind_invalid
@ kind_invalid
sycl::_V1::detail::kernel_param_desc_t::offset
int offset
Definition: kernel_desc.hpp:58
sycl::_V1::detail::kernel_param_kind_t::kind_pointer
@ kind_pointer
sycl::_V1::detail::SpecConstantInfo::getName
static constexpr const char * getName()
Definition: kernel_desc.hpp:63
sycl::_V1::detail::kernel_param_kind_t::kind_specialization_constants_buffer
@ kind_specialization_constants_buffer
sycl::_V1::detail::KernelInfo::getColumnNumber
static constexpr unsigned getColumnNumber()
Definition: kernel_desc.hpp:88
sycl::_V1::detail::kernel_param_kind_t::kind_sampler
@ kind_sampler
sycl::_V1::detail::KernelInfo::getParamDesc
static const kernel_param_desc_t & getParamDesc(int)
Definition: kernel_desc.hpp:79
sycl::_V1::detail::KernelInfo::getFileName
static constexpr const char * getFileName()
Definition: kernel_desc.hpp:85
sycl::_V1::detail::KernelInfo
Definition: kernel_desc.hpp:77
sycl::_V1::detail::KernelInfo::getKernelSize
static constexpr int64_t getKernelSize()
Definition: kernel_desc.hpp:89
sycl::_V1::detail::get_spec_constant_symbolic_ID_impl
const char * get_spec_constant_symbolic_ID_impl()
sycl::_V1::detail::KernelInfo::getNumParams
static constexpr unsigned getNumParams()
Definition: kernel_desc.hpp:78