Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

Commit 10ca23a

Browse files
jtmott-inteljtmott-intel
authored
Apply changes from intel/llvm #1963
1 parent 808b228 commit 10ca23a

File tree

3 files changed

+57
-47
lines changed

3 files changed

+57
-47
lines changed

‎SYCL/Basic/group-algorithm/exclusive_scan.cpp

Lines changed: 17 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717
using namespace sycl;
1818
using namespace sycl::ONEAPI;
1919

20-
template <class BinaryOperation, int TestNumber>
20+
template <class SpecializationKernelName, int TestNumber>
2121
class exclusive_scan_kernel;
2222

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

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

138138
#if __cplusplus >= 201402L
139-
test(q, input, output, plus<>(), 0);
140-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
141-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
139+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
140+
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
141+
std::numeric_limits<int>::max());
142+
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
143+
std::numeric_limits<int>::lowest());
142144
#endif
143-
test(q, input, output, plus<int>(), 0);
144-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
145-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
145+
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output, plus<int>(), 0);
146+
test<class KernelName_UXdGbr>(q, input, output, minimum<int>(),
147+
std::numeric_limits<int>::max());
148+
test<class KernelName_saYaodNyJknrPW>(q, input, output, maximum<int>(),
149+
std::numeric_limits<int>::lowest());
146150

147151
std::cout << "Test passed." << std::endl;
148152
}

‎SYCL/Basic/group-algorithm/inclusive_scan.cpp

Lines changed: 17 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717
using namespace sycl;
1818
using namespace sycl::ONEAPI;
1919

20-
template <class BinaryOperation, int TestNumber>
20+
template <class SpecializationKernelName, int TestNumber>
2121
class inclusive_scan_kernel;
2222

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

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

138138
#if __cplusplus >= 201402L
139-
test(q, input, output, plus<>(), 0);
140-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
141-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
139+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
140+
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
141+
std::numeric_limits<int>::max());
142+
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
143+
std::numeric_limits<int>::lowest());
142144
#endif
143-
test(q, input, output, plus<int>(), 0);
144-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
145-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
145+
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
146+
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
147+
std::numeric_limits<int>::max());
148+
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
149+
std::numeric_limits<int>::lowest());
146150

147151
std::cout << "Test passed." << std::endl;
148152
}

‎SYCL/Basic/group-algorithm/reduce.cpp

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

19-
template <class BinaryOperation>
20-
class reduce_kernel;
21-
22-
template <typename InputContainer, typename OutputContainer,
23-
class BinaryOperation>
19+
template <typename SpecializationKernelName, typename InputContainer,
20+
typename OutputContainer, class BinaryOperation>
2421
void test(queue q, InputContainer input, OutputContainer output,
2522
BinaryOperation binary_op,
2623
typename OutputContainer::value_type identity) {
2724
typedef typename InputContainer::value_type InputT;
2825
typedef typename OutputContainer::value_type OutputT;
29-
typedef class reduce_kernel<BinaryOperation> kernel_name;
3026
OutputT init = 42;
3127
size_t N = input.size();
3228
size_t G = 16;
@@ -37,15 +33,17 @@ void test(queue q, InputContainer input, OutputContainer output,
3733
q.submit([&](handler &cgh) {
3834
auto in = in_buf.template get_access<access::mode::read>(cgh);
3935
auto out = out_buf.template get_access<access::mode::discard_write>(cgh);
40-
cgh.parallel_for<kernel_name>(nd_range<1>(G, G), [=](nd_item<1> it) {
41-
group<1> g = it.get_group();
42-
int lid = it.get_local_id(0);
43-
out[0] = reduce(g, in[lid], binary_op);
44-
out[1] = reduce(g, in[lid], init, binary_op);
45-
out[2] = reduce(g, in.get_pointer(), in.get_pointer() + N, binary_op);
46-
out[3] =
47-
reduce(g, in.get_pointer(), in.get_pointer() + N, init, binary_op);
48-
});
36+
cgh.parallel_for<SpecializationKernelName>(
37+
nd_range<1>(G, G), [=](nd_item<1> it) {
38+
group<1> g = it.get_group();
39+
int lid = it.get_local_id(0);
40+
out[0] = reduce(g, in[lid], binary_op);
41+
out[1] = reduce(g, in[lid], init, binary_op);
42+
out[2] =
43+
reduce(g, in.get_pointer(), in.get_pointer() + N, binary_op);
44+
out[3] = reduce(g, in.get_pointer(), in.get_pointer() + N, init,
45+
binary_op);
46+
});
4947
});
5048
}
5149
// std::reduce is not implemented yet, so use std::accumulate instead
@@ -74,13 +72,17 @@ int main() {
7472
std::fill(output.begin(), output.end(), 0);
7573

7674
#if __cplusplus >= 201402L
77-
test(q, input, output, plus<>(), 0);
78-
test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
79-
test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
75+
test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
76+
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
77+
std::numeric_limits<int>::max());
78+
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
79+
std::numeric_limits<int>::lowest());
8080
#endif
81-
test(q, input, output, plus<int>(), 0);
82-
test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
83-
test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
81+
test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
82+
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
83+
std::numeric_limits<int>::max());
84+
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
85+
std::numeric_limits<int>::lowest());
8486

8587
std::cout << "Test passed." << std::endl;
8688
}

0 commit comments

Comments
 (0)