Skip to content

Commit 837f781

Browse files
FznamznonFznamznon
committed
Add test checking that spec ids work with device globals
1 parent 2538396 commit 837f781

File tree

5 files changed

+289
-4
lines changed

5 files changed

+289
-4
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -5076,9 +5076,6 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
50765076
Visited.insert(VD);
50775077
std::string TopShim = EmitSpecIdShims(OS, ShimCounter, Policy, VD);
50785078
if (Util::isSyclDeviceGlobalType(VD->getType())) {
5079-
if (!DeviceGlobalsEmitted)
5080-
OS << "#include <CL/sycl/detail/device_global_map.hpp>\n";
5081-
50825079
DeviceGlobalsEmitted = true;
50835080
DeviceGlobOS << "device_global_map::add(";
50845081
} else {
@@ -5127,6 +5124,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
51275124
OS << "#include <CL/sycl/detail/spec_const_integration.hpp>\n";
51285125

51295126
if (DeviceGlobalsEmitted) {
5127+
OS << "#include <CL/sycl/detail/device_global_map.hpp>\n";
51305128
DeviceGlobOS.flush();
51315129
OS << "namespace sycl::detail {\n";
51325130
OS << "namespace {\n";

‎clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,11 @@ __SYCL_INLINE_NAMESPACE(cl) {
66
namespace sycl {
77
namespace detail {
88

9+
#if __cplusplus >= 201703L
10+
template <auto &SpecName> const char *get_spec_constant_symbolic_ID_impl();
11+
template <auto &SpecName> const char *get_spec_constant_symbolic_ID();
12+
#endif
13+
914
#ifndef __SYCL_DEVICE_ONLY__
1015
#define _Bool bool
1116
#endif

‎clang/test/CodeGenSYCL/Inputs/CL/sycl/detail/spec_const_integration.hpp

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,15 @@
1+
__SYCL_INLINE_NAMESPACE(cl) {
2+
namespace sycl {
3+
namespace detail {
4+
5+
#if __cplusplus >= 201703L
6+
// Translates SYCL 2020 `specialization_id` to a unique symbolic identifier
7+
// which is used internally by the toolchain
8+
template <auto &SpecName> const char *get_spec_constant_symbolic_ID() {
9+
return get_spec_constant_symbolic_ID_impl<SpecName>();
10+
}
11+
#endif
12+
13+
} // namespace detail
14+
} // namespace sycl
15+
} // __SYCL_INLINE_NAMESPACE(cl)

‎clang/test/CodeGenSYCL/device_global_int_footer_header.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,6 @@ int main() {
3232
// CHECK-HEADER: } // __SYCL_INLINE_NAMESPACE(cl)
3333

3434
// CHECK-FOOTER: #include <CL/sycl/detail/defines_elementary.hpp>
35-
// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
3635

3736
// Shims go before the registration.
3837
// CHECK-FOOTER: namespace Foo {
@@ -52,6 +51,7 @@ int main() {
5251
// CHECK-FOOTER-NEXT: } // namespace __sycl_detail
5352
// CHECK-FOOTER-NEXT: } // namespace
5453

54+
// CHECK-FOOTER: #include <CL/sycl/detail/device_global_map.hpp>
5555
// CHECK-FOOTER: namespace sycl::detail {
5656
// CHECK-FOOTER-NEXT: namespace {
5757
// CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {

‎clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp

Lines changed: 267 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,267 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -fsycl-int-header=%t.header.h %s -emit-llvm -o %t.ll
2+
// RUN: FileCheck -input-file=%t.footer.h %s
3+
4+
// Try and compile all this stuff.
5+
// RUN: %clang_cc1 -fsycl-is-host -x c++ -std=c++17 -internal-isystem %S/Inputs -fsyntax-only -include %t.header.h -include %s %t.footer.h
6+
7+
// This test checks that integration footer is emitted correctly if both
8+
// spec constants and device globals are used.
9+
10+
#include "sycl.hpp"
11+
12+
using namespace cl;
13+
int main() {
14+
cl::sycl::kernel_single_task<class first_kernel>([]() {});
15+
}
16+
17+
// CHECK: #include <CL/sycl/detail/defines_elementary.hpp>
18+
constexpr sycl::specialization_id a{2};
19+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
20+
// CHECK-NEXT: namespace sycl {
21+
// CHECK-NEXT: namespace detail {
22+
// CHECK-NEXT: template<>
23+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::a>() {
24+
// CHECK-NEXT: return "____ZL1a";
25+
// CHECK-NEXT: }
26+
// CHECK-NEXT: } // namespace detail
27+
// CHECK-NEXT: } // namespace sycl
28+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
29+
sycl::ext::oneapi::device_global<int> b;
30+
31+
struct Wrapper {
32+
static constexpr sycl::specialization_id a{18};
33+
static sycl::ext::oneapi::device_global<float> b;
34+
};
35+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
36+
// CHECK-NEXT: namespace sycl {
37+
// CHECK-NEXT: namespace detail {
38+
// CHECK-NEXT: template<>
39+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::Wrapper::a>() {
40+
// CHECK-NEXT: return "_ZN7Wrapper1aE";
41+
// CHECK-NEXT: }
42+
// CHECK-NEXT: } // namespace detail
43+
// CHECK-NEXT: } // namespace sycl
44+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
45+
46+
template <typename T>
47+
struct TemplateWrapper {
48+
static constexpr sycl::specialization_id<T> a{18};
49+
static sycl::ext::oneapi::device_global<T> b;
50+
};
51+
52+
template class TemplateWrapper<float>;
53+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
54+
// CHECK-NEXT: namespace sycl {
55+
// CHECK-NEXT: namespace detail {
56+
// CHECK-NEXT: template<>
57+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::TemplateWrapper<float>::a>() {
58+
// CHECK-NEXT: return "_ZN15TemplateWrapperIfE1aE";
59+
// CHECK-NEXT: }
60+
// CHECK-NEXT: } // namespace detail
61+
// CHECK-NEXT: } // namespace sycl
62+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
63+
64+
namespace {
65+
constexpr sycl::specialization_id a{2};
66+
sycl::ext::oneapi::device_global<int> b;
67+
} // namespace
68+
69+
// CHECK: namespace {
70+
// CHECK-NEXT: namespace __sycl_detail {
71+
// CHECK-NEXT: static constexpr decltype(a) &__shim_[[SHIM0:[0-9]+]]() {
72+
// CHECK-NEXT: return a;
73+
// CHECK-NEXT: }
74+
// CHECK-NEXT: } // namespace __sycl_detail
75+
// CHECK-NEXT: } // namespace
76+
77+
// CHECK: __SYCL_INLINE_NAMESPACE(cl) {
78+
// CHECK-NEXT: namespace sycl {
79+
// CHECK-NEXT: namespace detail {
80+
// CHECK-NEXT: template<>
81+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::__sycl_detail::__shim_[[SHIM0]]()>() {
82+
// CHECK-NEXT: return "____ZN12_GLOBAL__N_11aE";
83+
// CHECK-NEXT: }
84+
// CHECK-NEXT: } // namespace detail
85+
// CHECK-NEXT: } // namespace sycl
86+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
87+
88+
// CHECK: namespace {
89+
// CHECK-NEXT: namespace __sycl_detail {
90+
// CHECK-NEXT: static constexpr decltype(b) &__shim_[[SHIM1:[0-9]+]]() {
91+
// CHECK-NEXT: return b;
92+
// CHECK-NEXT: }
93+
// CHECK-NEXT: } // namespace __sycl_detail
94+
// CHECK-NEXT: } // namespace
95+
96+
namespace outer {
97+
namespace {
98+
namespace inner {
99+
namespace {
100+
constexpr sycl::specialization_id a{2};
101+
// CHECK: namespace outer {
102+
// CHECK-NEXT: namespace {
103+
// CHECK-NEXT: namespace inner {
104+
// CHECK-NEXT: namespace {
105+
// CHECK-NEXT: namespace __sycl_detail {
106+
// CHECK-NEXT: static constexpr decltype(a) &__shim_[[SHIM2:[0-9]+]]() {
107+
// CHECK-NEXT: return a;
108+
// CHECK-NEXT: }
109+
// CHECK-NEXT: } // namespace __sycl_detail
110+
// CHECK-NEXT: } // namespace
111+
// CHECK-NEXT: } // namespace inner
112+
// CHECK-NEXT: } // namespace
113+
// CHECK-NEXT: } // namespace outer
114+
// CHECK-NEXT: namespace outer {
115+
// CHECK-NEXT: namespace {
116+
// CHECK-NEXT: namespace __sycl_detail {
117+
// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM2]]()) &__shim_[[SHIM3:[0-9]+]]() {
118+
// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM2]]();
119+
// CHECK-NEXT: }
120+
// CHECK-NEXT: } // namespace __sycl_detail
121+
// CHECK-NEXT: } // namespace
122+
// CHECK-NEXT: } // namespace outer
123+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
124+
// CHECK-NEXT: namespace sycl {
125+
// CHECK-NEXT: namespace detail {
126+
// CHECK-NEXT: template<>
127+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM3]]()>() {
128+
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_11aE";
129+
// CHECK-NEXT: }
130+
// CHECK-NEXT: } // namespace detail
131+
// CHECK-NEXT: } // namespace sycl
132+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
133+
sycl::ext::oneapi::device_global<int> b;
134+
// CHECK: namespace outer {
135+
// CHECK-NEXT: namespace {
136+
// CHECK-NEXT: namespace inner {
137+
// CHECK-NEXT: namespace {
138+
// CHECK-NEXT: namespace __sycl_detail {
139+
// CHECK-NEXT: static constexpr decltype(b) &__shim_[[SHIM4:[0-9]+]]() {
140+
// CHECK-NEXT: return b;
141+
// CHECK-NEXT: }
142+
// CHECK-NEXT: } // namespace __sycl_detail
143+
// CHECK-NEXT: } // namespace
144+
// CHECK-NEXT: } // namespace inner
145+
// CHECK-NEXT: } // namespace
146+
// CHECK-NEXT: } // namespace outer
147+
// CHECK-NEXT: namespace outer {
148+
// CHECK-NEXT: namespace {
149+
// CHECK-NEXT: namespace __sycl_detail {
150+
// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM4]]()) &__shim_[[SHIM5:[0-9]+]]() {
151+
// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM4]]();
152+
// CHECK-NEXT: }
153+
// CHECK-NEXT: } // namespace __sycl_detail
154+
// CHECK-NEXT: } // namespace
155+
// CHECK-NEXT: } // namespace outer
156+
struct Wrapper {
157+
static constexpr sycl::specialization_id a{18};
158+
static sycl::ext::oneapi::device_global<int> b;
159+
static sycl::ext::oneapi::device_global<float> c;
160+
};
161+
// CHECK: namespace outer {
162+
// CHECK-NEXT: namespace {
163+
// CHECK-NEXT: namespace inner {
164+
// CHECK-NEXT: namespace {
165+
// CHECK-NEXT: namespace __sycl_detail {
166+
// CHECK-NEXT: static constexpr decltype(Wrapper::a) &__shim_[[SHIM6:[0-9]+]]() {
167+
// CHECK-NEXT: return Wrapper::a;
168+
// CHECK-NEXT: }
169+
// CHECK-NEXT: } // namespace __sycl_detail
170+
// CHECK-NEXT: } // namespace
171+
// CHECK-NEXT: } // namespace inner
172+
// CHECK-NEXT: } // namespace
173+
// CHECK-NEXT: } // namespace outer
174+
// CHECK-NEXT: namespace outer {
175+
// CHECK-NEXT: namespace {
176+
// CHECK-NEXT: namespace __sycl_detail {
177+
// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM6]]()) &__shim_[[SHIM7:[0-9]+]]() {
178+
// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM6]]();
179+
// CHECK-NEXT: }
180+
// CHECK-NEXT: } // namespace __sycl_detail
181+
// CHECK-NEXT: } // namespace
182+
// CHECK-NEXT: } // namespace outer
183+
// CHECK-NEXT: __SYCL_INLINE_NAMESPACE(cl) {
184+
// CHECK-NEXT: namespace sycl {
185+
// CHECK-NEXT: namespace detail {
186+
// CHECK-NEXT: template<>
187+
// CHECK-NEXT: inline const char *get_spec_constant_symbolic_ID_impl<::outer::__sycl_detail::__shim_[[SHIM7]]()>() {
188+
// CHECK-NEXT: return "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1aE";
189+
// CHECK-NEXT: }
190+
// CHECK-NEXT: } // namespace detail
191+
// CHECK-NEXT: } // namespace sycl
192+
// CHECK-NEXT: } // __SYCL_INLINE_NAMESPACE(cl)
193+
194+
// CHECK-NEXT: namespace outer {
195+
// CHECK-NEXT: namespace {
196+
// CHECK-NEXT: namespace inner {
197+
// CHECK-NEXT: namespace {
198+
// CHECK-NEXT: namespace __sycl_detail {
199+
// CHECK-NEXT: static constexpr decltype(Wrapper::b) &__shim_[[SHIM8:[0-9]+]]() {
200+
// CHECK-NEXT: return Wrapper::b;
201+
// CHECK-NEXT: }
202+
// CHECK-NEXT: } // namespace __sycl_detail
203+
// CHECK-NEXT: } // namespace
204+
// CHECK-NEXT: } // namespace inner
205+
// CHECK-NEXT: } // namespace
206+
// CHECK-NEXT: } // namespace outer
207+
// CHECK-NEXT: namespace outer {
208+
// CHECK-NEXT: namespace {
209+
// CHECK-NEXT: namespace __sycl_detail {
210+
// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM8]]()) &__shim_[[SHIM9:[0-9]+]]() {
211+
// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM8]]();
212+
// CHECK-NEXT: }
213+
// CHECK-NEXT: } // namespace __sycl_detail
214+
// CHECK-NEXT: } // namespace
215+
// CHECK-NEXT: } // namespace outer
216+
// CHECK-NEXT: namespace outer {
217+
// CHECK-NEXT: namespace {
218+
// CHECK-NEXT: namespace inner {
219+
// CHECK-NEXT: namespace {
220+
// CHECK-NEXT: namespace __sycl_detail {
221+
// CHECK-NEXT: static constexpr decltype(Wrapper::c) &__shim_[[SHIM10:[0-9]+]]() {
222+
// CHECK-NEXT: return Wrapper::c;
223+
// CHECK-NEXT: }
224+
// CHECK-NEXT: } // namespace __sycl_detail
225+
// CHECK-NEXT: } // namespace
226+
// CHECK-NEXT: } // namespace inner
227+
// CHECK-NEXT: } // namespace
228+
// CHECK-NEXT: } // namespace outer
229+
// CHECK-NEXT: namespace outer {
230+
// CHECK-NEXT: namespace {
231+
// CHECK-NEXT: namespace __sycl_detail {
232+
// CHECK-NEXT: static constexpr decltype(inner::__sycl_detail::__shim_[[SHIM10]]()) &__shim_[[SHIM11:[0-9]+]]() {
233+
// CHECK-NEXT: return inner::__sycl_detail::__shim_[[SHIM10]]();
234+
// CHECK-NEXT: }
235+
// CHECK-NEXT: } // namespace __sycl_detail
236+
// CHECK-NEXT: } // namespace
237+
// CHECK-NEXT: } // namespace outer
238+
239+
// FIXME: Shims don't work with templated wrapper classes for some reason
240+
// template <typename T>
241+
// struct TemplateWrapper {
242+
// static constexpr sycl::specialization_id<T> a{18};
243+
// static sycl::ext::oneapi::device_global<T> b;
244+
// };
245+
//
246+
// template class TemplateWrapper<float>;
247+
248+
}
249+
}
250+
}
251+
}
252+
253+
// CHECK: #include <CL/sycl/detail/spec_const_integration.hpp>
254+
// CHECK-NEXT: #include <CL/sycl/detail/device_global_map.hpp>
255+
// CHECK-NEXT: namespace sycl::detail {
256+
// CHECK-NEXT: namespace {
257+
// CHECK-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
258+
// CHECK-NEXT: device_global_map::add((void *)&::b, "_Z1b");
259+
// CHECK-NEXT: device_global_map::add((void *)&::Wrapper::b, "_ZN7Wrapper1bE");
260+
// CHECK-NEXT: device_global_map::add((void *)&::TemplateWrapper<float>::b, "_ZN15TemplateWrapperIfE1bE");
261+
// CHECK-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "____ZN12_GLOBAL__N_11bE");
262+
// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM5]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_11bE");
263+
// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM9]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1bE");
264+
// CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM11]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1cE");
265+
// CHECK-NEXT: }
266+
// CHECK-NEXT: } // namespace (unnamed)
267+
// CHECK-NEXT: } // namespace sycl::detail

0 commit comments

Comments
 (0)