Skip to content

Commit f0d1636

Browse files
Fznamznonvladimirlaz
Fznamznon
authored and
vladimirlaz
committed
[SYCL] Switch to use one definition of SYCL runtime classes in SYCL tests
Signed-off-by: Vladimir Lazarev <[email protected]> Signed-off-by: Mariya Podchishchaeva <[email protected]>
1 parent ff0f21d commit f0d1636

File tree

3 files changed

+197
-248
lines changed

3 files changed

+197
-248
lines changed

‎clang/test/CodeGenSYCL/int_header1.cpp

Lines changed: 193 additions & 121 deletions
Original file line numberDiff line numberDiff line change
@@ -1,135 +1,207 @@
1-
// RUN: %clang --sycl -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv
1+
// RUN: %clang -I %S/Inputs --sycl -Xclang -fsycl-int-header=%t.h %s -c -o kernel.spv
2+
// RUN: %clang -I %S/Inputs --sycl -Xclang -fsycl-int-header=%t.h %s -c -o kernel.spv
23
// RUN: FileCheck -input-file=%t.h %s
3-
// XFAIL: *
4-
//
5-
// CHECK: class first_kernel;
6-
// CHECK-NEXT: template <typename T> class second_kernel;
7-
// CHECK-NEXT: struct X;
8-
// CHECK-NEXT: template <typename T> struct point ;
9-
// CHECK-NEXT: template <int a, typename T1, typename T2> class third_kernel;
10-
//
11-
// CHECK: #include <CL/sycl/detail/kernel_desc.hpp>
12-
//
13-
// CHECK: static constexpr
14-
// CHECK-NEXT: const char* const kernel_names[] = {
15-
// CHECK-NEXT: "first_kernel",
16-
// CHECK-NEXT: "second_namespace::second_kernel<char>",
17-
// CHECK-NEXT: "third_kernel<1, int, point< X> >"
18-
// CHECK-NEXT: };
19-
//
20-
// CHECK: const kernel_param_desc_t kernel_signatures[] = {
21-
// CHECK-NEXT: //--- first_kernel
22-
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
23-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2014, 4 },
24-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 5 },
25-
// CHECK-EMPTY:
26-
// CHECK-NEXT: //--- second_namespace::second_kernel<char>
27-
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
28-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
29-
// CHECK-EMPTY:
30-
// CHECK-NEXT: //--- third_kernel<1, int, point< X> >
31-
// CHECK-NEXT: { kernel_param_kind_t::kind_scalar, 4, 0 },
32-
// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 2016, 4 },
33-
// CHECK-EMPTY:
34-
// CHECK-NEXT: };
35-
//
36-
// CHECK: template <class KernelNameType> struct KernelInfo;
37-
// CHECK: template <> struct KernelInfo<class first_kernel> {
38-
// CHECK: template <> struct KernelInfo<class second_namespace::second_kernel<char>> {
39-
// CHECK: template <> struct KernelInfo<class third_kernel<1, int, struct point<struct X> >> {
40-
41-
namespace cl {
42-
namespace sycl {
43-
namespace access {
44-
45-
enum class target {
46-
global_buffer = 2014,
47-
constant_buffer,
48-
local,
49-
image,
50-
host_buffer,
51-
host_image,
52-
image_array
53-
};
544

55-
enum class mode {
56-
read = 1024,
57-
write,
58-
read_write,
59-
discard_write,
60-
discard_read_write,
61-
atomic
62-
};
5+
// CHECK:template <> struct KernelInfo<class KernelName> {
6+
// CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
7+
// CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
8+
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::nm2::KernelName0>> {
9+
// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::KernelName1>> {
10+
// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::nm2::KernelName0>> {
11+
// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::KernelName1>> {
12+
// CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
13+
// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
14+
// CHECK:template <> struct KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>> {
6315

64-
enum class placeholder { false_t,
65-
true_t };
16+
// This test checks if the SYCL device compiler is able to generate correct
17+
// integration header when the kernel name class is expressed in different
18+
// forms.
6619

67-
enum class address_space : int {
68-
private_space = 0,
69-
global_space,
70-
constant_space,
71-
local_space
72-
};
73-
} // namespace access
74-
template <typename dataT, int dimensions, access::mode accessmode,
75-
access::target accessTarget = access::target::global_buffer,
76-
access::placeholder isPlaceholder = access::placeholder::false_t>
77-
class accessor {
78-
79-
public:
80-
void use(void) const {}
81-
};
82-
} // namespace sycl
83-
} // namespace cl
20+
#include "sycl.hpp"
8421

8522
template <typename KernelName, typename KernelType>
8623
__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
8724
kernelFunc();
8825
}
89-
struct x {};
90-
template <typename T>
91-
struct point {};
92-
namespace second_namespace {
93-
template <typename T>
94-
class second_kernel;
26+
27+
namespace nm1 {
28+
namespace nm2 {
29+
class C {};
30+
class KernelName0 : public C {};
31+
} // namespace nm2
32+
33+
class KernelName1;
34+
35+
template <typename T> class KernelName3;
36+
template <typename T> class KernelName4;
37+
38+
template <> class KernelName3<nm1::nm2::KernelName0>;
39+
template <> class KernelName3<KernelName1>;
40+
41+
template <> class KernelName4<nm1::nm2::KernelName0> {};
42+
template <> class KernelName4<KernelName1> {};
43+
44+
} // namespace nm1
45+
46+
namespace {
47+
class ClassInAnonNS;
48+
template <typename T> class TmplClassInAnonNS;
9549
}
9650

97-
template <int a, typename T1, typename T2>
98-
class third_kernel;
51+
struct MyWrapper {
52+
class KN101 {};
9953

100-
int main() {
54+
int test() {
55+
56+
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc;
57+
58+
// Acronyms used to designate a test combination:
59+
// Declaration levels: 'T'-translation unit, 'L'-local scope,
60+
// 'C'-containing class, 'P'-"in place", '-'-N/A
61+
// Class definition: 'I'-incomplete (not defined), 'D' - defined,
62+
// '-'-N/A
63+
// Test combination positional parameters:
64+
// 0: Kernel class declaration level
65+
// 1: Kernel class definition
66+
// 2: Declaration level of the template argument class of the kernel class
67+
// 3: Definition of the template argument class of the kernel class
68+
69+
// PI--
70+
// traditional in-place incomplete type
71+
kernel_single_task<class KernelName>([=]() { acc.use(); });
72+
73+
// TD--
74+
// a class completely defined within a namespace at
75+
// translation unit scope
76+
kernel_single_task<nm1::nm2::KernelName0>([=]() { acc.use(); });
77+
78+
#ifdef LI__
79+
// TODO unexpected compilation error when host code + integration header
80+
// is compiled LI-- kernel name is an incomplete class forward-declared in
81+
// local scope
82+
class KernelName2;
83+
kernel_single_task<KernelName2>([=]() { acc.use(); });
84+
#endif
85+
86+
#ifdef LD__
87+
// Expected compilation error.
88+
// LD--
89+
// kernel name is a class defined in local scope
90+
class KernelName2a {};
91+
kernel_single_task<KernelName2a>([=]() { acc.use(); });
92+
#endif
10193

102-
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> acc1;
103-
cl::sycl::accessor<float, 2, cl::sycl::access::mode::write,
104-
cl::sycl::access::target::local,
105-
cl::sycl::access::placeholder::true_t>
106-
acc2;
107-
int i = 13;
108-
// TODO: Uncomemnt when structures in kernel arguments are correctly processed
109-
// by SYCL compiler
110-
/* struct {
111-
char c;
112-
int i;
113-
} test_s;
114-
test_s.c = 14;*/
115-
kernel_single_task<class first_kernel>([=]() {
116-
if (i == 13 /*&& test_s.c == 14*/) {
117-
118-
acc1.use();
119-
acc2.use();
120-
}
121-
});
122-
123-
kernel_single_task<class second_namespace::second_kernel<char>>([=]() {
124-
if (i == 13) {
125-
acc2.use();
126-
}
127-
});
128-
kernel_single_task<class third_kernel<1, int,point<struct X>>>([=]() {
129-
if (i == 13) {
130-
acc2.use();
131-
}
132-
});
133-
134-
return 0;
94+
// TI--
95+
// an incomplete class forward-declared in a namespace at
96+
// translation unit scope
97+
kernel_single_task<nm1::KernelName1>([=]() { acc.use(); });
98+
99+
// TITD
100+
// an incomplete template specialization class with defined class as
101+
// argument declared in a namespace at translation unit scope
102+
kernel_single_task<nm1::KernelName3<nm1::nm2::KernelName0>>(
103+
[=]() { acc.use(); });
104+
105+
// TITI
106+
// an incomplete template specialization class with incomplete class as
107+
// argument forward-declared in a namespace at translation unit scope
108+
kernel_single_task<nm1::KernelName3<nm1::KernelName1>>(
109+
[=]() { acc.use(); });
110+
111+
// TDTD
112+
// a defined template specialization class with defined class as argument
113+
// declared in a namespace at translation unit scope
114+
kernel_single_task<nm1::KernelName4<nm1::nm2::KernelName0>>(
115+
[=]() { acc.use(); });
116+
117+
// TDTI
118+
// a defined template specialization class with incomplete class as
119+
// argument forward-declared in a namespace at translation unit scope
120+
kernel_single_task<nm1::KernelName4<nm1::KernelName1>>(
121+
[=]() { acc.use(); });
122+
123+
// TIPI
124+
// an incomplete template specialization class with incomplete class as
125+
// argument forward-declared "in-place"
126+
kernel_single_task<nm1::KernelName3<class KernelName5>>(
127+
[=]() { acc.use(); });
128+
129+
#ifdef TILI
130+
// Expected compilation error
131+
// TILI
132+
// an incomplete template specialization class with incomplete class as
133+
// argument forward-declared locally
134+
class KernelName6;
135+
kernel_single_task<nm1::KernelName3<KernelName6>>(
136+
[=]() { acc.use(); });
137+
#endif
138+
139+
// TDPI
140+
// a defined template specialization class with incomplete class as
141+
// argument forward-declared "in-place"
142+
kernel_single_task<nm1::KernelName4<class KernelName7>>(
143+
[=]() { acc.use(); });
144+
145+
#ifdef TDLI
146+
// TODO unexpected compilation error when host code + integration header
147+
// is compiled TDLI a defined template specialization class with
148+
// incomplete class as argument forward-declared locally
149+
class KernelName6a;
150+
kernel_single_task<nm1::KernelName4<KernelName6a>>(
151+
[=]() { acc.use(); });
152+
#endif
153+
154+
#ifdef TDLD
155+
// Expected compilation error
156+
// TDLD
157+
// a defined template specialization class with a class as argument
158+
// defined locally
159+
class KernelName9 {};
160+
kernel_single_task<nm1::KernelName4<KernelName9>>([=]() { acc.use(); });
161+
#endif
162+
163+
#ifdef TICD
164+
// Expected compilation error
165+
// TICD
166+
// an incomplete template specialization class with a defined class as
167+
// argument declared in the containing class
168+
kernel_single_task<nm1::KernelName3<KN100>>([=]() { acc.use(); });
169+
#endif
170+
171+
#ifdef TICI
172+
// Expected compilation error
173+
// TICI
174+
// an incomplete template specialization class with an incomplete class as
175+
// argument declared in the containing class
176+
kernel_single_task<nm1::KernelName3<KN101>>([=]() { acc.use(); });
177+
#endif
178+
179+
// kernel name type is a templated class, both the top-level class and the
180+
// template argument are declared in the anonymous namespace
181+
kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(
182+
[=]() { acc.use(); });
183+
184+
return 0;
185+
}
186+
};
187+
188+
#ifndef __SYCL_DEVICE_ONLY__
189+
using namespace cl::sycl::detail;
190+
#endif // __SYCL_DEVICE_ONLY__
191+
192+
int main() {
193+
MyWrapper w;
194+
int a = w.test();
195+
#ifndef __SYCL_DEVICE_ONLY__
196+
KernelInfo<class KernelName>::getName();
197+
KernelInfo<class nm1::nm2::KernelName0>::getName();
198+
KernelInfo<class nm1::KernelName1>::getName();
199+
KernelInfo<class nm1::KernelName3<nm1::nm2::KernelName0>>::getName();
200+
KernelInfo<class nm1::KernelName3<class nm1::KernelName1>>::getName();
201+
KernelInfo<class nm1::KernelName4<nm1::nm2::KernelName0>>::getName();
202+
KernelInfo<class nm1::KernelName4<class nm1::KernelName1>>::getName();
203+
KernelInfo<class nm1::KernelName3<class KernelName5>>::getName();
204+
KernelInfo<class nm1::KernelName4<class KernelName7>>::getName();
205+
KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>>::getName();
206+
#endif //__SYCL_DEVICE_ONLY__
135207
}

0 commit comments

Comments
 (0)