Skip to content

Commit 2e16c13

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web' (#2)
2 parents 3425023 + ea73219 commit 2e16c13

11 files changed

+374
-344
lines changed

‎clang/test/SemaSYCL/built-in-type-kernel-arg.cpp

+32-24
Original file line numberDiff line numberDiff line change
@@ -1,14 +1,11 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -ast-dump %s | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s
22

33
// This test checks that compiler generates correct initialization for arguments
44
// that have struct or built-in type inside the OpenCL kernel
55

6-
#include "Inputs/sycl.hpp"
6+
#include "sycl.hpp"
77

8-
template <typename name, typename Func>
9-
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
10-
kernelFunc();
11-
}
8+
sycl::queue deviceQueue;
129

1310
struct test_struct {
1411
int data;
@@ -18,10 +15,12 @@ struct test_struct {
1815
};
1916

2017
void test(const int some_const) {
21-
kernel<class kernel_const>(
22-
[=]() {
23-
int a = some_const;
24-
});
18+
deviceQueue.submit([&](sycl::handler &h) {
19+
h.single_task<class kernel_const>(
20+
[=]() {
21+
int a = some_const;
22+
});
23+
});
2524
}
2625

2726
int main() {
@@ -31,20 +30,29 @@ int main() {
3130
int *ptr_array[2];
3231
test_struct s;
3332
s.data = data;
34-
kernel<class kernel_int>(
35-
[=]() {
36-
int kernel_data = data;
37-
});
38-
kernel<class kernel_struct>(
39-
[=]() {
40-
test_struct k_s;
41-
k_s = s;
42-
});
43-
kernel<class kernel_pointer>(
44-
[=]() {
45-
new_data_addr[0] = data_addr[0];
46-
int *local = ptr_array[1];
47-
});
33+
34+
deviceQueue.submit([&](sycl::handler &h) {
35+
h.single_task<class kernel_int>(
36+
[=]() {
37+
int kernel_data = data;
38+
});
39+
});
40+
41+
deviceQueue.submit([&](sycl::handler &h) {
42+
h.single_task<class kernel_struct>(
43+
[=]() {
44+
test_struct k_s;
45+
k_s = s;
46+
});
47+
});
48+
49+
deviceQueue.submit([&](sycl::handler &h) {
50+
h.single_task<class kernel_pointer>(
51+
[=]() {
52+
new_data_addr[0] = data_addr[0];
53+
int *local = ptr_array[1];
54+
});
55+
});
4856

4957
const int some_const = 10;
5058
test(some_const);

‎clang/test/SemaSYCL/call-to-undefined-function.cpp

+77-76
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,8 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -Wno-sycl-2017-compat -verify -fsyntax-only %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
2+
3+
#include "sycl.hpp"
4+
5+
sycl::queue deviceQueue;
26

37
void defined() {
48
}
@@ -8,11 +12,6 @@ void undefined();
812

913
SYCL_EXTERNAL void undefinedExternal();
1014

11-
template <typename name, typename Func>
12-
__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
13-
kernelFunc();
14-
}
15-
1615
template <typename T>
1716
void definedTpl() {
1817
}
@@ -95,76 +94,78 @@ int main() {
9594
// No problems in host code
9695
undefined();
9796

98-
kernel_single_task<class CallToUndefinedFnTester>([]() {
99-
// expected-note@-1 {{called by 'operator()'}}
100-
// expected-note@-2 {{called by 'operator()'}}
101-
102-
// simple functions
103-
defined();
104-
undefinedExternal();
105-
undefined();
106-
// expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
107-
108-
// templated functions
109-
definedTpl<int>();
110-
undefinedExternalTpl<int>();
111-
undefinedTpl<int>();
112-
// expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
113-
114-
// partially specialized template function
115-
definedPartialTpl<int, false>();
116-
definedPartialTpl<int, true>();
117-
definedPartialTpl<char, false>();
118-
definedPartialTpl<char, true>();
119-
120-
// template class with specialization
121-
{
122-
Tpl<int, false> tpl;
123-
tpl.defined();
124-
}
125-
126-
{
127-
Tpl<int, true> tpl;
128-
tpl.defined();
129-
}
130-
131-
// template class with template method, both have specializations.
132-
{
133-
TplWithTplMethod<int, false> tpl;
134-
tpl.defined<char, false>();
135-
tpl.defined<char, true>();
136-
tpl.defined<int, false>();
137-
tpl.defined<int, true>();
138-
}
139-
140-
{
141-
TplWithTplMethod<int, true> tpl;
142-
tpl.defined<char, false>();
143-
tpl.defined<char, true>();
144-
tpl.defined<int, false>();
145-
tpl.defined<int, true>();
146-
}
147-
148-
{
149-
TplWithTplMethod2<int, false> tpl;
150-
tpl.defined<char, false>();
151-
tpl.defined<char, true>();
152-
tpl.defined<int, false>();
153-
tpl.defined<int, true>();
154-
}
155-
156-
{
157-
TplWithTplMethod2<int, true> tpl;
158-
tpl.defined<char, false>();
159-
tpl.defined<char, true>();
160-
tpl.defined<int, false>();
161-
tpl.defined<int, true>();
162-
}
163-
164-
// forward-declared function
165-
useFwDeclFn();
166-
forwardDeclFn();
167-
forwardDeclFn2();
97+
deviceQueue.submit([&](sycl::handler &h) {
98+
h.single_task<class CallToUndefinedFnTester>([]() {
99+
// expected-note@-1 {{called by 'operator()'}}
100+
// expected-note@-2 {{called by 'operator()'}}
101+
102+
// simple functions
103+
defined();
104+
undefinedExternal();
105+
undefined();
106+
// expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
107+
108+
// templated functions
109+
definedTpl<int>();
110+
undefinedExternalTpl<int>();
111+
undefinedTpl<int>();
112+
// expected-error@-1 {{SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute}}
113+
114+
// partially specialized template function
115+
definedPartialTpl<int, false>();
116+
definedPartialTpl<int, true>();
117+
definedPartialTpl<char, false>();
118+
definedPartialTpl<char, true>();
119+
120+
// template class with specialization
121+
{
122+
Tpl<int, false> tpl;
123+
tpl.defined();
124+
}
125+
126+
{
127+
Tpl<int, true> tpl;
128+
tpl.defined();
129+
}
130+
131+
// template class with template method, both have specializations.
132+
{
133+
TplWithTplMethod<int, false> tpl;
134+
tpl.defined<char, false>();
135+
tpl.defined<char, true>();
136+
tpl.defined<int, false>();
137+
tpl.defined<int, true>();
138+
}
139+
140+
{
141+
TplWithTplMethod<int, true> tpl;
142+
tpl.defined<char, false>();
143+
tpl.defined<char, true>();
144+
tpl.defined<int, false>();
145+
tpl.defined<int, true>();
146+
}
147+
148+
{
149+
TplWithTplMethod2<int, false> tpl;
150+
tpl.defined<char, false>();
151+
tpl.defined<char, true>();
152+
tpl.defined<int, false>();
153+
tpl.defined<int, true>();
154+
}
155+
156+
{
157+
TplWithTplMethod2<int, true> tpl;
158+
tpl.defined<char, false>();
159+
tpl.defined<char, true>();
160+
tpl.defined<int, false>();
161+
tpl.defined<int, true>();
162+
}
163+
164+
// forward-declared function
165+
useFwDeclFn();
166+
forwardDeclFn();
167+
forwardDeclFn2();
168+
});
168169
});
169170
}
170171

‎clang/test/SemaSYCL/deferred-diagnostics-aux-builtin.cpp

+11-15
Original file line numberDiff line numberDiff line change
@@ -1,16 +1,8 @@
1-
// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl -fsycl-is-device -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat -verify -fsyntax-only %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-unknown-linux-gnu -Wno-sycl-2017-compat -verify -fsyntax-only %s
22

3-
inline namespace cl {
4-
namespace sycl {
3+
#include "sycl.hpp"
54

6-
template <typename name, typename Func>
7-
__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
8-
// expected-note@+1 {{called by 'kernel_single_task<AName, (lambda}}
9-
kernelFunc();
10-
}
11-
12-
} // namespace sycl
13-
} // namespace cl
5+
sycl::queue deviceQueue;
146

157
int main(int argc, char **argv) {
168
//_mm_prefetch is an x86-64 specific builtin where the second integer parameter is required to be a constant
@@ -19,9 +11,13 @@ int main(int argc, char **argv) {
1911

2012
_mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}}
2113

22-
cl::sycl::kernel_single_task<class AName>([]() {
23-
_mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}}
24-
_mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}}
14+
deviceQueue.submit([&](sycl::handler &h) {
15+
// expected-note@Inputs/sycl.hpp:212 {{called by 'kernel_single_task<AName, (lambda}}
16+
h.single_task<class AName>([]() {
17+
_mm_prefetch("test", 4); // expected-error {{builtin is not supported on this target}}
18+
_mm_prefetch("test", 8); // expected-error {{argument value 8 is outside the valid range [0, 7]}} expected-error {{builtin is not supported on this target}}
19+
});
2520
});
21+
2622
return 0;
27-
}
23+
}

0 commit comments

Comments
 (0)