Skip to content

Commit 96d2e17

Browse files
schittirschittir
authored
[SYCL] Add clang support for code_location in KernelInfo (#5335)
Add code_location support in clang This patch adds four methods to KernelInfo class to return filename, kernel name, line number and column number. These methods return meaningful information if NDEBUG is not defined, else, they emit empty string or 0.
1 parent 68efbec commit 96d2e17

File tree

3 files changed

+362
-0
lines changed

3 files changed

+362
-0
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4721,6 +4721,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
47214721

47224722
for (const KernelDesc &K : KernelDescs) {
47234723
const size_t N = K.Params.size();
4724+
PresumedLoc PLoc = S.Context.getSourceManager().getPresumedLoc(
4725+
S.Context.getSourceManager()
4726+
.getExpansionRange(K.KernelLocation)
4727+
.getEnd());
47244728
if (K.IsUnnamedKernel) {
47254729
O << "template <> struct KernelInfoData<";
47264730
OutputStableNameInChars(O, K.StableName);
@@ -4744,6 +4748,44 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
47444748
O << " __SYCL_DLL_LOCAL\n";
47454749
O << " static constexpr bool isESIMD() { return " << K.IsESIMDKernel
47464750
<< "; }\n";
4751+
O << " __SYCL_DLL_LOCAL\n";
4752+
O << " static constexpr const char* getFileName() {\n";
4753+
O << "#ifndef NDEBUG\n";
4754+
O << " return \""
4755+
<< std::string(PLoc.getFilename())
4756+
.substr(std::string(PLoc.getFilename()).find_last_of("/\\") + 1);
4757+
O << "\";\n";
4758+
O << "#else\n";
4759+
O << " return \"\";\n";
4760+
O << "#endif\n";
4761+
O << " }\n";
4762+
O << " __SYCL_DLL_LOCAL\n";
4763+
O << " static constexpr const char* getFunctionName() {\n";
4764+
O << "#ifndef NDEBUG\n";
4765+
O << " return \"";
4766+
SYCLKernelNameTypePrinter Printer(O, Policy);
4767+
Printer.Visit(K.NameType);
4768+
O << "\";\n";
4769+
O << "#else\n";
4770+
O << " return \"\";\n";
4771+
O << "#endif\n";
4772+
O << " }\n";
4773+
O << " __SYCL_DLL_LOCAL\n";
4774+
O << " static constexpr unsigned getLineNumber() {\n";
4775+
O << "#ifndef NDEBUG\n";
4776+
O << " return " << PLoc.getLine() << ";\n";
4777+
O << "#else\n";
4778+
O << " return 0;\n";
4779+
O << "#endif\n";
4780+
O << " }\n";
4781+
O << " __SYCL_DLL_LOCAL\n";
4782+
O << " static constexpr unsigned getColumnNumber() {\n";
4783+
O << "#ifndef NDEBUG\n";
4784+
O << " return " << PLoc.getColumn() << ";\n";
4785+
O << "#else\n";
4786+
O << " return 0;\n";
4787+
O << "#endif\n";
4788+
O << " }\n";
47474789
O << "};\n";
47484790
CurStart += N;
47494791
}

‎clang/test/CodeGenSYCL/code_location.cpp

Lines changed: 308 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,308 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
2+
// RUN: FileCheck -input-file=%t.h %s
3+
// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2020 %s | FileCheck -input-file=%t.h %s
4+
5+
#include "Inputs/sycl.hpp"
6+
7+
// Check that meaningful information is returned when NDEBUG is not defined
8+
// and empty strings and 0s are emitted when it is.
9+
int test1() {
10+
cl::sycl::queue q;
11+
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
12+
q.submit([&](cl::sycl::handler &h) { h.single_task<class KernelName>([]() {}); });
13+
return 0;
14+
}
15+
// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', '5', 't', 'e', 's', 't', '1', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '2', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
16+
// CHECK: static constexpr const char* getFileName() {
17+
// CHECK: #ifndef NDEBUG
18+
// CHECK: return "code_location.cpp";
19+
// CHECK: #else
20+
// CHECK: return "";
21+
// CHECK: #endif
22+
// CHECK: }
23+
// CHECK: static constexpr const char* getFunctionName() {
24+
// CHECK: #ifndef NDEBUG
25+
// CHECK: return "";
26+
// CHECK: #else
27+
// CHECK: return "";
28+
// CHECK: #endif
29+
// CHECK: }
30+
// CHECK: static constexpr unsigned getLineNumber() {
31+
// CHECK: #ifndef NDEBUG
32+
// CHECK: return 11;
33+
// CHECK: #else
34+
// CHECK: return 0;
35+
// CHECK: #endif
36+
// CHECK: }
37+
// CHECK: static constexpr unsigned getColumnNumber() {
38+
// CHECK: #ifndef NDEBUG
39+
// CHECK: return 54;
40+
// CHECK: #else
41+
// CHECK: return 0;
42+
// CHECK: #endif
43+
// CHECK: }
44+
// CHECK: };
45+
46+
// CHECK: template <> struct KernelInfo<KernelName> {
47+
// CHECK: static constexpr const char* getFileName() {
48+
// CHECK: #ifndef NDEBUG
49+
// CHECK: return "code_location.cpp";
50+
// CHECK: #else
51+
// CHECK: return "";
52+
// CHECK: #endif
53+
// CHECK: }
54+
// CHECK: __SYCL_DLL_LOCAL
55+
// CHECK: static constexpr const char* getFunctionName() {
56+
// CHECK: #ifndef NDEBUG
57+
// CHECK: return "KernelName";
58+
// CHECK: #else
59+
// CHECK: return "";
60+
// CHECK: #endif
61+
// CHECK: }
62+
// CHECK: __SYCL_DLL_LOCAL
63+
// CHECK: static constexpr unsigned getLineNumber() {
64+
// CHECK: #ifndef NDEBUG
65+
// CHECK: return 12;
66+
// CHECK: #else
67+
// CHECK: return 0;
68+
// CHECK: #endif
69+
// CHECK: }
70+
// CHECK: __SYCL_DLL_LOCAL
71+
// CHECK: static constexpr unsigned getColumnNumber() {
72+
// CHECK: #ifndef NDEBUG
73+
// CHECK: return 72;
74+
// CHECK: #else
75+
// CHECK: return 0;
76+
// CHECK: #endif
77+
// CHECK: }
78+
// CHECK: };
79+
80+
// Check that the right name and location is returned when
81+
// lambda and kernel name are defined on different lines
82+
class KernelName2;
83+
int test2() {
84+
cl::sycl::queue q;
85+
q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName2>(
86+
[] { int i = 2; }); });
87+
return 0;
88+
}
89+
// CHECK: template <> struct KernelInfo<::KernelName2> {
90+
// CHECK: static constexpr const char* getFileName() {
91+
// CHECK: #ifndef NDEBUG
92+
// CHECK: return "code_location.cpp";
93+
// CHECK: #else
94+
// CHECK: return "";
95+
// CHECK: #endif
96+
// CHECK: }
97+
// CHECK: static constexpr const char* getFunctionName() {
98+
// CHECK: #ifndef NDEBUG
99+
// CHECK: return "::KernelName2";
100+
// CHECK: #else
101+
// CHECK: return "";
102+
// CHECK: #endif
103+
// CHECK: }
104+
// CHECK: static constexpr unsigned getLineNumber() {
105+
// CHECK: #ifndef NDEBUG
106+
// CHECK: return 86;
107+
// CHECK: #else
108+
// CHECK: return 0;
109+
// CHECK: #endif
110+
// CHECK: }
111+
// CHECK: static constexpr unsigned getColumnNumber() {
112+
// CHECK: #ifndef NDEBUG
113+
// CHECK: return 44;
114+
// CHECK: #else
115+
// CHECK: return 0;
116+
// CHECK: #endif
117+
// CHECK: }
118+
// CHECK: };
119+
120+
// Check that fully qualified name is returned
121+
template <typename T> class KernelName3;
122+
int test3() {
123+
cl::sycl::queue q;
124+
q.submit([&](cl::sycl::handler &h) { h.single_task<KernelName3<KernelName2>>(
125+
[] { int i = 3; }); });
126+
return 0;
127+
}
128+
// CHECK: template <> struct KernelInfo<::KernelName3<::KernelName2>> {
129+
// CHECK: static constexpr const char* getFileName() {
130+
// CHECK: #ifndef NDEBUG
131+
// CHECK: return "code_location.cpp";
132+
// CHECK: #else
133+
// CHECK: return "";
134+
// CHECK: #endif
135+
// CHECK: }
136+
// CHECK: static constexpr const char* getFunctionName() {
137+
// CHECK: #ifndef NDEBUG
138+
// CHECK: return "::KernelName3<::KernelName2>";
139+
// CHECK: #else
140+
// CHECK: return "";
141+
// CHECK: #endif
142+
// CHECK: }
143+
// CHECK: static constexpr unsigned getLineNumber() {
144+
// CHECK: #ifndef NDEBUG
145+
// CHECK: return 125;
146+
// CHECK: #else
147+
// CHECK: return 0;
148+
// CHECK: #endif
149+
// CHECK: }
150+
// CHECK: static constexpr unsigned getColumnNumber() {
151+
// CHECK: #ifndef NDEBUG
152+
// CHECK: return 44;
153+
// CHECK: #else
154+
// CHECK: return 0;
155+
// CHECK: #endif
156+
// CHECK: }
157+
// CHECK: };
158+
159+
// Check that the location information returned is that of l4
160+
auto l4 = []() { return 4; };
161+
int test4() {
162+
cl::sycl::queue q;
163+
q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName4>(l4); });
164+
return 0;
165+
}
166+
// CHECK: template <> struct KernelInfo<KernelName4> {
167+
// CHECK: static constexpr const char* getFileName() {
168+
// CHECK: #ifndef NDEBUG
169+
// CHECK: return "code_location.cpp";
170+
// CHECK: #else
171+
// CHECK: return "";
172+
// CHECK: #endif
173+
// CHECK: }
174+
// CHECK: static constexpr const char* getFunctionName() {
175+
// CHECK: #ifndef NDEBUG
176+
// CHECK: return "KernelName4";
177+
// CHECK: #else
178+
// CHECK: return "";
179+
// CHECK: #endif
180+
// CHECK: }
181+
// CHECK: static constexpr unsigned getLineNumber() {
182+
// CHECK: #ifndef NDEBUG
183+
// CHECK: return 160;
184+
// CHECK: #else
185+
// CHECK: return 0;
186+
// CHECK: #endif
187+
// CHECK: }
188+
// CHECK: static constexpr unsigned getColumnNumber() {
189+
// CHECK: #ifndef NDEBUG
190+
// CHECK: return 11;
191+
// CHECK: #else
192+
// CHECK: return 0;
193+
// CHECK: #endif
194+
// CHECK: }
195+
// CHECK: };
196+
197+
// Check that fully qualified name is returned when unnamed lambda
198+
// kernel is enclosed in a namespace
199+
namespace NS {
200+
int test5() {
201+
cl::sycl::queue q;
202+
q.submit([=](cl::sycl::handler &h) { h.single_task([] {}); });
203+
q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName5>([] {}); });
204+
return 0;
205+
}
206+
} // namespace NS
207+
// CHECK: template <> struct KernelInfoData<'_', 'Z', 'T', 'S', 'Z', 'Z', 'N', '2', 'N', 'S', '5', 't', 'e', 's', 't', '5', 'E', 'v', 'E', 'N', 'K', 'U', 'l', 'R', 'N', '2', 'c', 'l', '4', 's', 'y', 'c', 'l', '7', 'h', 'a', 'n', 'd', 'l', 'e', 'r', 'E', 'E', '_', 'c', 'l', 'E', 'S', '3', '_', 'E', 'U', 'l', 'v', 'E', '_'> {
208+
// CHECK: static constexpr const char* getFileName() {
209+
// CHECK: #ifndef NDEBUG
210+
// CHECK: return "code_location.cpp";
211+
// CHECK: #else
212+
// CHECK: return "";
213+
// CHECK: #endif
214+
// CHECK: }
215+
// CHECK: static constexpr const char* getFunctionName() {
216+
// CHECK: #ifndef NDEBUG
217+
// CHECK: return "NS::";
218+
// CHECK: #else
219+
// CHECK: return "";
220+
// CHECK: #endif
221+
// CHECK: }
222+
// CHECK: static constexpr unsigned getLineNumber() {
223+
// CHECK: #ifndef NDEBUG
224+
// CHECK: return 202;
225+
// CHECK: #else
226+
// CHECK: return 0;
227+
// CHECK: #endif
228+
// CHECK: }
229+
// CHECK: static constexpr unsigned getColumnNumber() {
230+
// CHECK: #ifndef NDEBUG
231+
// CHECK: return 54;
232+
// CHECK: #else
233+
// CHECK: return 0;
234+
// CHECK: #endif
235+
// CHECK: }
236+
// CHECK: };
237+
// CHECK: template <> struct KernelInfo<NS::KernelName5> {
238+
// CHECK: static constexpr const char* getFileName() {
239+
// CHECK: #ifndef NDEBUG
240+
// CHECK: return "code_location.cpp";
241+
// CHECK: #else
242+
// CHECK: return "";
243+
// CHECK: #endif
244+
// CHECK: }
245+
// CHECK: static constexpr const char* getFunctionName() {
246+
// CHECK: #ifndef NDEBUG
247+
// CHECK: return "NS::KernelName5";
248+
// CHECK: #else
249+
// CHECK: return "";
250+
// CHECK: #endif
251+
// CHECK: }
252+
// CHECK: static constexpr unsigned getLineNumber() {
253+
// CHECK: #ifndef NDEBUG
254+
// CHECK: return 203;
255+
// CHECK: #else
256+
// CHECK: return 0;
257+
// CHECK: #endif
258+
// CHECK: }
259+
// CHECK: static constexpr unsigned getColumnNumber() {
260+
// CHECK: #ifndef NDEBUG
261+
// CHECK: return 73;
262+
// CHECK: #else
263+
// CHECK: return 0;
264+
// CHECK: #endif
265+
// CHECK: }
266+
// CHECK: };
267+
268+
// Check that the location information returned is that of the Functor
269+
struct Functor {
270+
void operator()() const {
271+
}
272+
};
273+
int test6() {
274+
Functor F;
275+
cl::sycl::queue q;
276+
q.submit([=](cl::sycl::handler &h) { h.single_task<class KernelName6>(F); });
277+
return 0;
278+
}
279+
// CHECK: template <> struct KernelInfo<KernelName6> {
280+
// CHECK: static constexpr const char* getFileName() {
281+
// CHECK: #ifndef NDEBUG
282+
// CHECK: return "code_location.cpp";
283+
// CHECK: #else
284+
// CHECK: return "";
285+
// CHECK: #endif
286+
// CHECK: }
287+
// CHECK: static constexpr const char* getFunctionName() {
288+
// CHECK: #ifndef NDEBUG
289+
// CHECK: return "KernelName6";
290+
// CHECK: #else
291+
// CHECK: return "";
292+
// CHECK: #endif
293+
// CHECK: }
294+
// CHECK: static constexpr unsigned getLineNumber() {
295+
// CHECK: #ifndef NDEBUG
296+
// CHECK: return 269;
297+
// CHECK: #else
298+
// CHECK: return 0;
299+
// CHECK: #endif
300+
// CHECK: }
301+
// CHECK: static constexpr unsigned getColumnNumber() {
302+
// CHECK: #ifndef NDEBUG
303+
// CHECK: return 8;
304+
// CHECK: #else
305+
// CHECK: return 0;
306+
// CHECK: #endif
307+
// CHECK: }
308+
// CHECK: };

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

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -75,6 +75,10 @@ template <class KernelNameType> struct KernelInfo {
7575
}
7676
static constexpr const char *getName() { return ""; }
7777
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; }
7882
};
7983
#else
8084
template <char...> struct KernelInfoData {
@@ -85,6 +89,10 @@ template <char...> struct KernelInfoData {
8589
}
8690
static constexpr const char *getName() { return ""; }
8791
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; }
8896
};
8997

9098
// C++14 like index_sequence and make_index_sequence
@@ -123,6 +131,10 @@ template <class KernelNameType> struct KernelInfo {
123131
}
124132
static constexpr const char *getName() { return SubKernelInfo::getName(); }
125133
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; }
126138
};
127139
#endif //__SYCL_UNNAMED_LAMBDA__
128140

0 commit comments

Comments
 (0)