Skip to content

Commit d5b9385

Browse files
aelovikov-intelaelovikov-intel
committed
[NFCI][SYCL] More use of CompileTimeKernelInfoTy/DeviceKernelInfo
Additional cleanups enabled by intel#19117.
1 parent d058a13 commit d5b9385

25 files changed

+302
-320
lines changed

‎sycl/include/sycl/detail/compile_time_kernel_info.hpp‎

Lines changed: 90 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -10,41 +10,117 @@
1010
#include <sycl/detail/kernel_desc.hpp>
1111
#include <sycl/detail/string_view.hpp>
1212

13+
#include <cassert>
14+
1315
namespace sycl {
1416
inline namespace _V1 {
1517
namespace detail {
16-
inline namespace compile_time_kernel_info_v1 {
1718

19+
template <typename KernelNameType>
20+
constexpr kernel_param_desc_t getKernelParamDesc(int Idx) {
21+
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
22+
kernel_param_desc_t ParamDesc;
23+
ParamDesc.kind =
24+
__builtin_sycl_kernel_param_kind(KernelIdentity<KernelNameType>(), Idx);
25+
ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor
26+
? __builtin_sycl_kernel_param_access_target(
27+
KernelIdentity<KernelNameType>(), Idx)
28+
: __builtin_sycl_kernel_param_size(
29+
KernelIdentity<KernelNameType>(), Idx);
30+
ParamDesc.offset =
31+
__builtin_sycl_kernel_param_offset(KernelIdentity<KernelNameType>(), Idx);
32+
return ParamDesc;
33+
#else
34+
return KernelInfo<KernelNameType>::getParamDesc(Idx);
35+
#endif
36+
}
37+
38+
inline namespace compile_time_kernel_info_v1 {
1839
// This is being passed across ABI boundary, so we don't use std::string_view,
1940
// at least for as long as we support user apps built with GNU libstdc++'s
2041
// pre-C++11 ABI.
2142
struct CompileTimeKernelInfoTy {
22-
detail::string_view Name;
43+
detail::string_view Name{};
2344
unsigned NumParams = 0;
2445
bool IsESIMD = false;
46+
// TODO: Can we just have code_location here?
2547
detail::string_view FileName{};
2648
detail::string_view FunctionName{};
2749
unsigned LineNumber = 0;
2850
unsigned ColumnNumber = 0;
2951
int64_t KernelSize = 0;
3052
using ParamDescGetterT = kernel_param_desc_t (*)(int);
3153
ParamDescGetterT ParamDescGetter = nullptr;
32-
bool HasSpecialCaptures = true;
54+
55+
bool HasSpecialCaptures = [this]() constexpr {
56+
// No-compile time info for the kernel (i.e., kernel_bundle/interop/etc.),
57+
// be conservative:
58+
if (NumParams == 0)
59+
return true;
60+
61+
bool FoundSpecialCapture = false;
62+
for (unsigned I = 0; I < NumParams; ++I) {
63+
auto ParamDesc = ParamDescGetter(I);
64+
bool IsSpecialCapture =
65+
(ParamDesc.kind != kernel_param_kind_t::kind_std_layout &&
66+
ParamDesc.kind != kernel_param_kind_t::kind_pointer);
67+
FoundSpecialCapture |= IsSpecialCapture;
68+
}
69+
return FoundSpecialCapture;
70+
}();
71+
72+
void refine(const CompileTimeKernelInfoTy &Other) {
73+
auto RefineField = [](auto &Self, const auto &Other, const auto &Default) {
74+
if (Self == Default) {
75+
Self = Other;
76+
} else if (Other != Default) {
77+
assert(Self == Other);
78+
}
79+
};
80+
81+
constexpr CompileTimeKernelInfoTy Empty;
82+
RefineField(Name, static_cast<std::string_view>(Other.Name),
83+
static_cast<std::string_view>(Empty.Name));
84+
RefineField(NumParams, Other.NumParams, Empty.NumParams);
85+
RefineField(IsESIMD, Other.IsESIMD, Empty.IsESIMD);
86+
RefineField(FileName, static_cast<std::string_view>(Other.FileName),
87+
static_cast<std::string_view>(Empty.FileName));
88+
RefineField(FunctionName, static_cast<std::string_view>(Other.FunctionName),
89+
static_cast<std::string_view>(Empty.FunctionName));
90+
RefineField(LineNumber, Other.LineNumber, Empty.LineNumber);
91+
RefineField(ColumnNumber, Other.ColumnNumber, Empty.ColumnNumber);
92+
RefineField(KernelSize, Other.KernelSize, Empty.KernelSize);
93+
RefineField(ParamDescGetter, Other.ParamDescGetter, Empty.ParamDescGetter);
94+
RefineField(HasSpecialCaptures, Other.HasSpecialCaptures,
95+
Empty.HasSpecialCaptures);
96+
}
3397
};
3498

3599
template <class Kernel>
36100
inline constexpr CompileTimeKernelInfoTy CompileTimeKernelInfo{
37-
std::string_view(getKernelName<Kernel>()),
38-
getKernelNumParams<Kernel>(),
39-
isKernelESIMD<Kernel>(),
40-
std::string_view(getKernelFileName<Kernel>()),
41-
std::string_view(getKernelFunctionName<Kernel>()),
42-
getKernelLineNumber<Kernel>(),
43-
getKernelColumnNumber<Kernel>(),
44-
getKernelSize<Kernel>(),
45-
&getKernelParamDesc<Kernel>,
46-
hasSpecialCaptures<Kernel>()};
47-
101+
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
102+
__builtin_sycl_kernel_name(KernelIdentity<Kernel>()),
103+
__builtin_sycl_kernel_param_count(KernelIdentity<Kernel>()),
104+
false /*IsESIMD*/, // TODO needs a builtin counterpart
105+
__builtin_sycl_kernel_file_name(KernelIdentity<Kernel>()),
106+
__builtin_sycl_kernel_function_name(KernelIdentity<Kernel>()),
107+
__builtin_sycl_kernel_line_number(KernelIdentity<Kernel>()),
108+
__builtin_sycl_kernel_column_number(KernelIdentity<Kernel>()),
109+
// TODO needs a builtin counterpart, but is currently only used for checking
110+
// cases with external host compiler, which use integration headers.
111+
0 /* KernelSize */, &getKernelParamDesc<Kernel>
112+
#else
113+
detail::string_view{KernelInfo<Kernel>::getName()},
114+
KernelInfo<Kernel>::getNumParams(), KernelInfo<Kernel>::isESIMD(),
115+
detail::string_view{KernelInfo<Kernel>::getFileName()},
116+
detail::string_view{KernelInfo<Kernel>::getFunctionName()},
117+
KernelInfo<Kernel>::getLineNumber(), KernelInfo<Kernel>::getColumnNumber(),
118+
KernelInfo<Kernel>::getKernelSize(),
119+
// Can't use KernelInfo::getParamDesc due to different return type (const
120+
// ref vs. by val):
121+
&getKernelParamDesc<Kernel>
122+
#endif
123+
};
48124
} // namespace compile_time_kernel_info_v1
49125
} // namespace detail
50126
} // namespace _V1

‎sycl/include/sycl/detail/kernel_desc.hpp‎

Lines changed: 0 additions & 90 deletions
Original file line numberDiff line numberDiff line change
@@ -187,96 +187,6 @@ template <typename KNT> struct KernelIdentity {
187187
using type = KNT;
188188
};
189189

190-
template <typename KernelNameType> constexpr unsigned getKernelNumParams() {
191-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
192-
return __builtin_sycl_kernel_param_count(KernelIdentity<KernelNameType>());
193-
#else
194-
return KernelInfo<KernelNameType>::getNumParams();
195-
#endif
196-
}
197-
198-
template <typename KernelNameType>
199-
constexpr kernel_param_desc_t getKernelParamDesc(int Idx) {
200-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
201-
kernel_param_desc_t ParamDesc;
202-
ParamDesc.kind =
203-
__builtin_sycl_kernel_param_kind(KernelIdentity<KernelNameType>(), Idx);
204-
ParamDesc.info = ParamDesc.kind == kernel_param_kind_t::kind_accessor
205-
? __builtin_sycl_kernel_param_access_target(
206-
KernelIdentity<KernelNameType>(), Idx)
207-
: __builtin_sycl_kernel_param_size(
208-
KernelIdentity<KernelNameType>(), Idx);
209-
ParamDesc.offset =
210-
__builtin_sycl_kernel_param_offset(KernelIdentity<KernelNameType>(), Idx);
211-
return ParamDesc;
212-
#else
213-
return KernelInfo<KernelNameType>::getParamDesc(Idx);
214-
#endif
215-
}
216-
217-
template <typename KernelNameType> constexpr const char *getKernelName() {
218-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
219-
return __builtin_sycl_kernel_name(KernelIdentity<KernelNameType>());
220-
#else
221-
return KernelInfo<KernelNameType>::getName();
222-
#endif
223-
}
224-
225-
template <typename KernelNameType> constexpr bool isKernelESIMD() {
226-
// TODO Needs a builtin counterpart
227-
return KernelInfo<KernelNameType>::isESIMD();
228-
}
229-
230-
template <typename KernelNameType> constexpr const char *getKernelFileName() {
231-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
232-
return __builtin_sycl_kernel_file_name(KernelIdentity<KernelNameType>());
233-
#else
234-
return KernelInfo<KernelNameType>::getFileName();
235-
#endif
236-
}
237-
238-
template <typename KernelNameType>
239-
constexpr const char *getKernelFunctionName() {
240-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
241-
return __builtin_sycl_kernel_function_name(KernelIdentity<KernelNameType>());
242-
#else
243-
return KernelInfo<KernelNameType>::getFunctionName();
244-
#endif
245-
}
246-
247-
template <typename KernelNameType> constexpr unsigned getKernelLineNumber() {
248-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
249-
return __builtin_sycl_kernel_line_number(KernelIdentity<KernelNameType>());
250-
#else
251-
return KernelInfo<KernelNameType>::getLineNumber();
252-
#endif
253-
}
254-
255-
template <typename KernelNameType> constexpr unsigned getKernelColumnNumber() {
256-
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
257-
return __builtin_sycl_kernel_column_number(KernelIdentity<KernelNameType>());
258-
#else
259-
return KernelInfo<KernelNameType>::getColumnNumber();
260-
#endif
261-
}
262-
263-
template <typename KernelNameType> constexpr int64_t getKernelSize() {
264-
// TODO needs a builtin counterpart, but is currently only used for checking
265-
// cases with external host compiler, which use integration headers.
266-
return KernelInfo<KernelNameType>::getKernelSize();
267-
}
268-
269-
template <typename KernelNameType> constexpr bool hasSpecialCaptures() {
270-
bool FoundSpecialCapture = false;
271-
for (unsigned I = 0; I < getKernelNumParams<KernelNameType>(); ++I) {
272-
auto ParamDesc = getKernelParamDesc<KernelNameType>(I);
273-
bool IsSpecialCapture =
274-
(ParamDesc.kind != kernel_param_kind_t::kind_std_layout &&
275-
ParamDesc.kind != kernel_param_kind_t::kind_pointer);
276-
FoundSpecialCapture |= IsSpecialCapture;
277-
}
278-
return FoundSpecialCapture;
279-
}
280190
} // namespace detail
281191
} // namespace _V1
282192
} // namespace sycl

‎sycl/include/sycl/detail/kernel_launch_helper.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
#pragma once
1010

1111
#include <sycl/detail/cg_types.hpp>
12+
#include <sycl/detail/compile_time_kernel_info.hpp>
1213
#include <sycl/detail/helpers.hpp>
1314
#include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
1415
#include <sycl/ext/intel/experimental/kernel_execution_properties.hpp>
@@ -261,7 +262,8 @@ struct KernelLaunchPropertyWrapper {
261262
if constexpr (ext::oneapi::experimental::detail::
262263
HasKernelPropertiesGetMethod<const KernelType &>::value) {
263264

264-
h->template processProperties<detail::isKernelESIMD<KernelName>()>(
265+
h->template processProperties<
266+
detail::CompileTimeKernelInfo<KernelName>.IsESIMD>(
265267
KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
266268
}
267269
#endif

‎sycl/include/sycl/detail/string_view.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -38,6 +38,8 @@ class string_view {
3838
#endif
3939
{
4040
}
41+
explicit constexpr string_view(const char *str) noexcept
42+
: string_view(std::string_view{str}) {}
4143
string_view(const sycl::detail::string &strn) noexcept
4244
: str(strn.c_str())
4345
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES

0 commit comments

Comments
 (0)