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 
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  static constexpr const char *getFileName() { return ""; }
79  static constexpr const char *getFunctionName() { return ""; }
80  static constexpr unsigned getLineNumber() { return 0; }
81  static constexpr unsigned getColumnNumber() { return 0; }
82 };
83 #else
84 template <char...> struct KernelInfoData {
85  static constexpr unsigned getNumParams() { return 0; }
86  static const kernel_param_desc_t &getParamDesc(int Idx) {
87  static kernel_param_desc_t Dummy;
88  return Dummy;
89  }
90  static constexpr const char *getName() { return ""; }
91  static constexpr bool isESIMD() { return 0; }
92  static constexpr const char *getFileName() { return ""; }
93  static constexpr const char *getFunctionName() { return ""; }
94  static constexpr unsigned getLineNumber() { return 0; }
95  static constexpr unsigned getColumnNumber() { return 0; }
96 };
97 
98 // C++14 like index_sequence and make_index_sequence
99 // not needed C++14 members (value_type, size) not implemented
100 template <class T, T...> struct integer_sequence {};
101 template <unsigned long long... I>
102 using index_sequence = integer_sequence<unsigned long long, I...>;
103 template <unsigned long long N>
104 using make_index_sequence =
105  __make_integer_seq<integer_sequence, unsigned long long, N>;
106 
107 template <typename T> struct KernelInfoImpl {
108 private:
109  static constexpr auto n = __builtin_sycl_unique_stable_name(T);
110  template <unsigned long long... I>
111  static KernelInfoData<n[I]...> impl(index_sequence<I...>) {
112  return {};
113  }
114 
115 public:
116  using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{}));
117 };
118 
119 // For named kernels, this structure is specialized in the integration header.
120 // For unnamed kernels, KernelInfoData is specialized in the integration header,
121 // and this picks it up via the KernelInfoImpl. For non-existent kernels, this
122 // will also pick up a KernelInfoData (as SubKernelInfo) via KernelInfoImpl, but
123 // it will instead get the unspecialized case, defined above.
124 template <class KernelNameType> struct KernelInfo {
125  using SubKernelInfo = typename KernelInfoImpl<KernelNameType>::type;
126  static constexpr unsigned getNumParams() {
127  return SubKernelInfo::getNumParams();
128  }
129  static const kernel_param_desc_t &getParamDesc(int Idx) {
130  return SubKernelInfo::getParamDesc(Idx);
131  }
132  static constexpr const char *getName() { return SubKernelInfo::getName(); }
133  static constexpr bool isESIMD() { return SubKernelInfo::isESIMD(); }
134  static constexpr const char *getFileName() { return ""; }
135  static constexpr const char *getFunctionName() { return ""; }
136  static constexpr unsigned getLineNumber() { return 0; }
137  static constexpr unsigned getColumnNumber() { return 0; }
138 };
139 #endif //__SYCL_UNNAMED_LAMBDA__
140 
141 } // namespace detail
142 } // namespace sycl
143 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::KernelInfo::getColumnNumber
static constexpr unsigned getColumnNumber()
Definition: kernel_desc.hpp:81
cl::sycl::detail::KernelInfo::getName
static constexpr const char * getName()
Definition: kernel_desc.hpp:76
T
defines_elementary.hpp
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
sycl
Definition: invoke_simd.hpp:68
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::KernelInfo::getFileName
static constexpr const char * getFileName()
Definition: kernel_desc.hpp:78
cl::sycl::detail::KernelInfo::getLineNumber
static constexpr unsigned getLineNumber()
Definition: kernel_desc.hpp:80
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::KernelInfo::getFunctionName
static constexpr const char * getFunctionName()
Definition: kernel_desc.hpp:79
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