Skip to content

Commit d2d55b0

Browse files
committed
Fixing tests, specifically tests with std kernels
1 parent 1d345ec commit d2d55b0

File tree

8 files changed

+101
-107
lines changed

8 files changed

+101
-107
lines changed

clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -DCHECK_ERROR -verify %s
2-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s
2+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s
33
// RUN: FileCheck -input-file=%t.h %s
44
//
55
// CHECK: #include <CL/sycl/detail/defines.hpp>

clang/test/SemaSYCL/unnamed-kernel.cpp

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -47,15 +47,15 @@ struct MyWrapper {
4747

4848
#ifndef __SYCL_UNNAMED_LAMBDA__
4949
// expected-error@+4 {{kernel needs to have a globally-visible name}}
50-
// expected-note@21 {{InvalidKernelName0 declared here}}
50+
// expected-note@22 {{InvalidKernelName0 declared here}}
5151
#endif
5252
q.submit([&](cl::sycl::handler &h) {
5353
h.single_task<InvalidKernelName0>([] {});
5454
});
5555

5656
#ifndef __SYCL_UNNAMED_LAMBDA__
5757
// expected-error@+4 {{kernel needs to have a globally-visible name}}
58-
// expected-note@22 {{InvalidKernelName3 declared here}}
58+
// expected-note@23 {{InvalidKernelName3 declared here}}
5959
#endif
6060
q.submit([&](cl::sycl::handler &h) {
6161
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
@@ -67,7 +67,7 @@ struct MyWrapper {
6767
});
6868

6969
#ifndef __SYCL_UNNAMED_LAMBDA__
70-
// expected-error@+3 {{kernel name is missing}}
70+
// expected-error@+3 {{kernel name cannot be a type in the "std" namespace}}
7171
#endif
7272
q.submit([&](cl::sycl::handler &h) {
7373
h.single_task<std::max_align_t>([] {});
@@ -76,7 +76,7 @@ struct MyWrapper {
7676
using InvalidAlias = InvalidKernelName4;
7777
#ifndef __SYCL_UNNAMED_LAMBDA__
7878
// expected-error@+4 {{kernel needs to have a globally-visible name}}
79-
// expected-note@23 {{InvalidKernelName4 declared here}}
79+
// expected-note@24 {{InvalidKernelName4 declared here}}
8080
#endif
8181
q.submit([&](cl::sycl::handler &h) {
8282
h.single_task<InvalidAlias>([] {});
@@ -85,7 +85,7 @@ struct MyWrapper {
8585
using InvalidAlias1 = InvalidKernelName5;
8686
#ifndef __SYCL_UNNAMED_LAMBDA__
8787
// expected-error@+4 {{kernel needs to have a globally-visible name}}
88-
// expected-note@24 {{InvalidKernelName5 declared here}}
88+
// expected-note@25 {{InvalidKernelName5 declared here}}
8989
#endif
9090
q.submit([&](cl::sycl::handler &h) {
9191
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});

sycl/test/group-algorithm/exclusive_scan.cpp

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616
using namespace sycl;
1717
using namespace sycl::intel;
1818

19-
template <class BinaryOperation, int TestNumber>
19+
template <class SpecializationKernelName, int TestNumber>
2020
class exclusive_scan_kernel;
2121

2222
// std::exclusive_scan isn't implemented yet, so use serial implementation
@@ -36,17 +36,17 @@ OutputIterator exclusive_scan(InputIterator first, InputIterator last,
3636
}
3737
} // namespace emu
3838

39-
template <typename InputContainer, typename OutputContainer,
40-
class BinaryOperation>
39+
template <typename SpecializationKernelName, typename InputContainer,
40+
typename OutputContainer, class BinaryOperation>
4141
void test(queue q, InputContainer input, OutputContainer output,
4242
BinaryOperation binary_op,
4343
typename OutputContainer::value_type identity) {
4444
typedef typename InputContainer::value_type InputT;
4545
typedef typename OutputContainer::value_type OutputT;
46-
typedef class exclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
47-
typedef class exclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
48-
typedef class exclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
49-
typedef class exclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
46+
typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
47+
typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
48+
typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
49+
typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
5050
OutputT init = 42;
5151
size_t N = input.size();
5252
size_t G = 16;
@@ -135,13 +135,13 @@ int main() {
135135
std::fill(output.begin(), output.end(), 0);
136136

137137
#if __cplusplus >= 201402L
138-
test(q, input, output, plus<>(), 0);
139-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
140-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
138+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
139+
test<class KernelNameMinimumV>(q, input, output, minimum<>(), std::numeric_limits<int>::max());
140+
test<class KernelNameMaximumV>(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
141141
#endif
142-
test(q, input, output, plus<int>(), 0);
143-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
144-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
142+
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
143+
test<class KernelNameMinimumI>(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
144+
test<class KernelNameMaximumI>(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
145145

146146
std::cout << "Test passed." << std::endl;
147147
}

sycl/test/group-algorithm/inclusive_scan.cpp

Lines changed: 13 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616
using namespace sycl;
1717
using namespace sycl::intel;
1818

19-
template <class BinaryOperation, int TestNumber>
19+
template <class SpecializationKernelName, int TestNumber>
2020
class inclusive_scan_kernel;
2121

2222
// std::inclusive_scan isn't implemented yet, so use serial implementation
@@ -36,17 +36,17 @@ OutputIterator inclusive_scan(InputIterator first, InputIterator last,
3636
}
3737
} // namespace emu
3838

39-
template <typename InputContainer, typename OutputContainer,
40-
class BinaryOperation>
39+
template <typename SpecializationKernelName, typename InputContainer,
40+
typename OutputContainer, class BinaryOperation>
4141
void test(queue q, InputContainer input, OutputContainer output,
4242
BinaryOperation binary_op,
4343
typename OutputContainer::value_type identity) {
4444
typedef typename InputContainer::value_type InputT;
4545
typedef typename OutputContainer::value_type OutputT;
46-
typedef class inclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
47-
typedef class inclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
48-
typedef class inclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
49-
typedef class inclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
46+
typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
47+
typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
48+
typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
49+
typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
5050
OutputT init = 42;
5151
size_t N = input.size();
5252
size_t G = 16;
@@ -135,13 +135,13 @@ int main() {
135135
std::fill(output.begin(), output.end(), 0);
136136

137137
#if __cplusplus >= 201402L
138-
test(q, input, output, plus<>(), 0);
139-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
140-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
138+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
139+
test<class KernelNameMinimumV>(q, input, output, minimum<>(), std::numeric_limits<int>::max());
140+
test<class KernelNameMaximumV>(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
141141
#endif
142-
test(q, input, output, plus<int>(), 0);
143-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
144-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
142+
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
143+
test<class KernelNameMinimumI>(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
144+
test<class KernelNameMaximumI>(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
145145

146146
std::cout << "Test passed." << std::endl;
147147
}

sycl/test/group-algorithm/reduce.cpp

Lines changed: 9 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -15,17 +15,13 @@
1515
using namespace sycl;
1616
using namespace sycl::intel;
1717

18-
template <class BinaryOperation>
19-
class reduce_kernel;
20-
21-
template <typename InputContainer, typename OutputContainer,
22-
class BinaryOperation>
18+
template <typename SpecializationKernelName, typename InputContainer,
19+
typename OutputContainer, class BinaryOperation>
2320
void test(queue q, InputContainer input, OutputContainer output,
2421
BinaryOperation binary_op,
2522
typename OutputContainer::value_type identity) {
2623
typedef typename InputContainer::value_type InputT;
2724
typedef typename OutputContainer::value_type OutputT;
28-
typedef class reduce_kernel<BinaryOperation> kernel_name;
2925
OutputT init = 42;
3026
size_t N = input.size();
3127
size_t G = 16;
@@ -36,7 +32,7 @@ void test(queue q, InputContainer input, OutputContainer output,
3632
q.submit([&](handler &cgh) {
3733
auto in = in_buf.template get_access<access::mode::read>(cgh);
3834
auto out = out_buf.template get_access<access::mode::discard_write>(cgh);
39-
cgh.parallel_for<kernel_name>(nd_range<1>(G, G), [=](nd_item<1> it) {
35+
cgh.parallel_for<SpecializationKernelName>(nd_range<1>(G, G), [=](nd_item<1> it) {
4036
group<1> g = it.get_group();
4137
int lid = it.get_local_id(0);
4238
out[0] = reduce(g, in[lid], binary_op);
@@ -73,13 +69,13 @@ int main() {
7369
std::fill(output.begin(), output.end(), 0);
7470

7571
#if __cplusplus >= 201402L
76-
test(q, input, output, plus<>(), 0);
77-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
78-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
72+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
73+
test<class KernelNameMinimumV>(q, input, output, minimum<>(), std::numeric_limits<int>::max());
74+
test<class KernelNameMaximumV>(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
7975
#endif
80-
test(q, input, output, plus<int>(), 0);
81-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
82-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
76+
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
77+
test<class KernelNameMinimumI>(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
78+
test<class KernelNameMaximumI>(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
8379

8480
std::cout << "Test passed." << std::endl;
8581
}

sycl/test/reduction/reduction_nd_ext_type.hpp

Lines changed: 10 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -8,10 +8,8 @@
88

99
using namespace cl::sycl;
1010

11-
template <typename T, int Dim, class BinaryOperation>
12-
class SomeClass;
13-
14-
template <typename T, int Dim, access::mode Mode, class BinaryOperation>
11+
template <typename SpecializationKernelName, typename T,
12+
int Dim, access::mode Mode, class BinaryOperation>
1513
void test(T Identity, size_t WGSize, size_t NWItems) {
1614
buffer<T, 1> InBuf(NWItems);
1715
buffer<T, 1> OutBuf(1);
@@ -35,7 +33,7 @@ void test(T Identity, size_t WGSize, size_t NWItems) {
3533
range<1> GlobalRange(NWItems);
3634
range<1> LocalRange(WGSize);
3735
nd_range<1> NDRange(GlobalRange, LocalRange);
38-
CGH.parallel_for<SomeClass<T, Dim, BinaryOperation>>(
36+
CGH.parallel_for<SpecializationKernelName>(
3937
NDRange, Redu, [=](nd_item<1> NDIt, auto &Sum) {
4038
Sum.combine(In[NDIt.get_global_linear_id()]);
4139
});
@@ -63,16 +61,16 @@ int runTests(const string_class &ExtensionName) {
6361
}
6462

6563
// Check some less standards WG sizes and corner cases first.
66-
test<T, 1, access::mode::read_write, std::multiplies<T>>(0, 4, 4);
67-
test<T, 0, access::mode::discard_write, intel::plus<T>>(0, 4, 64);
64+
test<class KernelName_oTh, T, 1, access::mode::read_write, std::multiplies<T>>(0, 4, 4);
65+
test<class KernelName_QUQnMARQT, T, 0, access::mode::discard_write, intel::plus<T>>(0, 4, 64);
6866

69-
test<T, 0, access::mode::read_write, intel::minimum<T>>(getMaximumFPValue<T>(), 7, 7);
70-
test<T, 1, access::mode::discard_write, intel::maximum<T>>(getMinimumFPValue<T>(), 7, 7 * 5);
67+
test<class KernelName_xGixNo, T, 0, access::mode::read_write, intel::minimum<T>>(getMaximumFPValue<T>(), 7, 7);
68+
test<class KernelName_qXNFw, T, 1, access::mode::discard_write, intel::maximum<T>>(getMinimumFPValue<T>(), 7, 7 * 5);
7169

7270
#if __cplusplus >= 201402L
73-
test<T, 1, access::mode::read_write, intel::plus<>>(1, 3, 3 * 5);
74-
test<T, 1, access::mode::discard_write, intel::minimum<>>(getMaximumFPValue<T>(), 3, 3);
75-
test<T, 0, access::mode::discard_write, intel::maximum<>>(getMinimumFPValue<T>(), 3, 3);
71+
test<class KernelName_lXdWtzANdDcvm, T, 1, access::mode::read_write, intel::plus<>>(1, 3, 3 * 5);
72+
test<class KernelName_FDQalsDxmbi, T, 1, access::mode::discard_write, intel::minimum<>>(getMaximumFPValue<T>(), 3, 3);
73+
test<class KernelName_TaNRRxDRXbzYrFImPYC, T, 0, access::mode::discard_write, intel::maximum<>>(getMinimumFPValue<T>(), 3, 3);
7674
#endif // __cplusplus >= 201402L
7775

7876
std::cout << "Test passed\n";

sycl/test/sub_group/reduce.cpp

Lines changed: 23 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -19,20 +19,20 @@
1919
#include "helper.hpp"
2020
#include <CL/sycl.hpp>
2121

22-
template <typename T, class BinaryOperation>
22+
template <typename ... Ts>
2323
class sycl_subgr;
2424

2525
using namespace cl::sycl;
2626

27-
template <typename T, class BinaryOperation>
27+
template <typename SpecializationKernelName, typename T, class BinaryOperation>
2828
void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
2929
size_t G = 240, size_t L = 60) {
3030
try {
3131
nd_range<1> NdRange(G, L);
3232
buffer<T> buf(G);
3333
Queue.submit([&](handler &cgh) {
3434
auto acc = buf.template get_access<access::mode::read_write>(cgh);
35-
cgh.parallel_for<sycl_subgr<T, BinaryOperation>>(
35+
cgh.parallel_for<SpecializationKernelName>(
3636
NdRange, [=](nd_item<1> NdItem) {
3737
intel::sub_group sg = NdItem.get_sub_group();
3838
if (skip_init) {
@@ -70,32 +70,32 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
7070
}
7171
}
7272

73-
template <typename T>
73+
template <typename SpecializationKernelName, typename T>
7474
void check(queue &Queue, size_t G = 240, size_t L = 60) {
7575
// limit data range for half to avoid rounding issues
7676
if (std::is_same<T, cl::sycl::half>::value) {
7777
G = 64;
7878
L = 32;
7979
}
8080

81-
check_op<T>(Queue, T(L), intel::plus<T>(), false, G, L);
82-
check_op<T>(Queue, T(0), intel::plus<T>(), true, G, L);
81+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_cNsJzXxSBQfEKY>, T>(Queue, T(L), intel::plus<T>(), false, G, L);
82+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_bWdCJaxe>, T>(Queue, T(0), intel::plus<T>(), true, G, L);
8383

84-
check_op<T>(Queue, T(0), intel::minimum<T>(), false, G, L);
85-
check_op<T>(Queue, T(G), intel::minimum<T>(), true, G, L);
84+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_wjspvpHJtI>, T>(Queue, T(0), intel::minimum<T>(), false, G, L);
85+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_BUioaQYxhjN>, T>(Queue, T(G), intel::minimum<T>(), true, G, L);
8686

87-
check_op<T>(Queue, T(G), intel::maximum<T>(), false, G, L);
88-
check_op<T>(Queue, T(0), intel::maximum<T>(), true, G, L);
87+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_bIHcoJBNpiB>, T>(Queue, T(G), intel::maximum<T>(), false, G, L);
88+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_bPPlfvdGShi>, T>(Queue, T(0), intel::maximum<T>(), true, G, L);
8989

9090
#if __cplusplus >= 201402L
91-
check_op<T>(Queue, T(L), intel::plus<>(), false, G, L);
92-
check_op<T>(Queue, T(0), intel::plus<>(), true, G, L);
91+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_fkOyLRYirfMnvBcnbRFy>, T>(Queue, T(L), intel::plus<>(), false, G, L);
92+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_zhzfRmSAFlswKWShyecv>, T>(Queue, T(0), intel::plus<>(), true, G, L);
9393

94-
check_op<T>(Queue, T(0), intel::minimum<>(), false, G, L);
95-
check_op<T>(Queue, T(G), intel::minimum<>(), true, G, L);
94+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_NaOzDnOmDPiDIXnXvaGy>, T>(Queue, T(0), intel::minimum<>(), false, G, L);
95+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_XXAfdcNmCNX>, T>(Queue, T(G), intel::minimum<>(), true, G, L);
9696

97-
check_op<T>(Queue, T(G), intel::maximum<>(), false, G, L);
98-
check_op<T>(Queue, T(0), intel::maximum<>(), true, G, L);
97+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_pLlvjjZsPv>, T>(Queue, T(G), intel::maximum<>(), false, G, L);
98+
check_op<sycl_subgr<SpecializationKernelName, class KernelName_BaCGaWDMFeMFqvotbk>, T>(Queue, T(0), intel::maximum<>(), true, G, L);
9999
#endif
100100
}
101101

@@ -106,19 +106,19 @@ int main() {
106106
return 0;
107107
}
108108

109-
check<int>(Queue);
110-
check<unsigned int>(Queue);
111-
check<long>(Queue);
112-
check<unsigned long>(Queue);
113-
check<float>(Queue);
109+
check<class KernelName_AJprOaCZgUmsYFRTTGNw, int>(Queue);
110+
check<class KernelName_ShKFIYTqaI, unsigned int>(Queue);
111+
check<class KernelName_TovsKTk, long>(Queue);
112+
check<class KernelName_JqbvoN, unsigned long>(Queue);
113+
check<class KernelName_mAWqKSWTT, float>(Queue);
114114
// reduce half type is not supported in OCL CPU RT
115115
#ifdef SG_GPU
116116
if (Queue.get_device().has_extension("cl_khr_fp16")) {
117-
check<cl::sycl::half>(Queue);
117+
check<class KernelName_CNwvDdLPNleaRvSuYr, cl::sycl::half>(Queue);
118118
}
119119
#endif
120120
if (Queue.get_device().has_extension("cl_khr_fp64")) {
121-
check<double>(Queue);
121+
check<class KernelName_bqHXlkbm, double>(Queue);
122122
}
123123
std::cout << "Test passed." << std::endl;
124124
return 0;

0 commit comments

Comments
 (0)