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

[SYCL] Specified ONEAPI aliases in functions and function objects #300

Merged
merged 9 commits into from
Jun 4, 2021
Merged

[SYCL] Specified ONEAPI aliases in functions and function objects #300

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
9 commits
Select commit Hold shift + click to select a range
591bd92
[SYCL] Added test for constexpr half operations
maximdimakov May 24, 2021
f549f76
Merge remote-tracking branch 'origin/intel' into intel
maximdimakov May 24, 2021
e442d4b
Merge remote-tracking branch 'origin/intel' into intel
maximdimakov May 27, 2021
0acb1ec
[SYCL] Specified ONEAPI namespace in functions
maximdimakov May 28, 2021
3369f9d
Reverted changes for half
maximdimakov May 28, 2021
31f0c2f
Clang-format fix
maximdimakov May 28, 2021
972f94e
Provided more ONEAPI aliases
maximdimakov Jun 3, 2021
a737723
Added explicit aliases for ONEAPI function objects
maximdimakov Jun 3, 2021
04a0be2
Clang-format fix
maximdimakov Jun 3, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 13 additions & 11 deletions SYCL/GroupAlgorithm/exclusive_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -140,24 +140,26 @@ int main() {
std::iota(input.begin(), input.end(), 0);
std::fill(output.begin(), output.end(), 0);

test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
test<class KernelNamePlusV>(q, input, output, ONEAPI::plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, ONEAPI::minimum<>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
test<class KernelNameMaximumV>(q, input, output, ONEAPI::maximum<>(),
std::numeric_limits<int>::lowest());

test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
test<class KernelNamePlusI>(q, input, output, ONEAPI::plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, ONEAPI::minimum<int>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
test<class KernelNameMaximumI>(q, input, output, ONEAPI::maximum<int>(),
std::numeric_limits<int>::lowest());

#ifdef SPIRV_1_3
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output, multiplies<int>(),
1);
test<class KernelName_UXdGbr>(q, input, output, bit_or<int>(), 0);
test<class KernelName_saYaodNyJknrPW>(q, input, output, bit_xor<int>(), 0);
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, bit_and<int>(), ~0);
test<class KernelName_VzAPutpBRRJrQPB>(q, input, output,
ONEAPI::multiplies<int>(), 1);
test<class KernelName_UXdGbr>(q, input, output, ONEAPI::bit_or<int>(), 0);
test<class KernelName_saYaodNyJknrPW>(q, input, output,
ONEAPI::bit_xor<int>(), 0);
test<class KernelName_GPcuAlvAOjrDyP>(q, input, output,
ONEAPI::bit_and<int>(), ~0);
#endif // SPIRV_1_3

std::cout << "Test passed." << std::endl;
Expand Down
24 changes: 13 additions & 11 deletions SYCL/GroupAlgorithm/inclusive_scan.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -140,25 +140,27 @@ int main() {
std::iota(input.begin(), input.end(), 0);
std::fill(output.begin(), output.end(), 0);

test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
test<class KernelNamePlusV>(q, input, output, ONEAPI::plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, ONEAPI::minimum<>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
test<class KernelNameMaximumV>(q, input, output, ONEAPI::maximum<>(),
std::numeric_limits<int>::lowest());

test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
test<class KernelNamePlusI>(q, input, output, ONEAPI::plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, ONEAPI::minimum<int>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
test<class KernelNameMaximumI>(q, input, output, ONEAPI::maximum<int>(),
std::numeric_limits<int>::lowest());

#ifdef SPIRV_1_3
test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
multiplies<int>(), 1);
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output, bit_or<int>(), 0);
test<class KernelName_yXIZfjwjxQGiPeQAnc>(q, input, output, bit_xor<int>(),
0);
test<class KernelName_xGnAnMYHvqekCk>(q, input, output, bit_and<int>(), ~0);
ONEAPI::multiplies<int>(), 1);
test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output,
ONEAPI::bit_or<int>(), 0);
test<class KernelName_yXIZfjwjxQGiPeQAnc>(q, input, output,
ONEAPI::bit_xor<int>(), 0);
test<class KernelName_xGnAnMYHvqekCk>(q, input, output,
ONEAPI::bit_and<int>(), ~0);
#endif // SPIRV_1_3

std::cout << "Test passed." << std::endl;
Expand Down
22 changes: 12 additions & 10 deletions SYCL/GroupAlgorithm/reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -76,24 +76,26 @@ int main() {
std::iota(input.begin(), input.end(), 0);
std::fill(output.begin(), output.end(), 0);

test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, minimum<>(),
test<class KernelNamePlusV>(q, input, output, ONEAPI::plus<>(), 0);
test<class KernelNameMinimumV>(q, input, output, ONEAPI::minimum<>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumV>(q, input, output, maximum<>(),
test<class KernelNameMaximumV>(q, input, output, ONEAPI::maximum<>(),
std::numeric_limits<int>::lowest());

test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
test<class KernelNamePlusI>(q, input, output, ONEAPI::plus<int>(), 0);
test<class KernelNameMinimumI>(q, input, output, ONEAPI::minimum<int>(),
std::numeric_limits<int>::max());
test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
test<class KernelNameMaximumI>(q, input, output, ONEAPI::maximum<int>(),
std::numeric_limits<int>::lowest());

#ifdef SPIRV_1_3
test<class KernelName_WonwuUVPUPOTKRKIBtT>(q, input, output,
multiplies<int>(), 1);
test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output, bit_or<int>(), 0);
test<class KernelName_eLSFt>(q, input, output, bit_xor<int>(), 0);
test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output, bit_and<int>(), ~0);
ONEAPI::multiplies<int>(), 1);
test<class KernelName_qYBaJDZTMGkdIwD>(q, input, output,
ONEAPI::bit_or<int>(), 0);
test<class KernelName_eLSFt>(q, input, output, ONEAPI::bit_xor<int>(), 0);
test<class KernelName_uFhJnxSVhNAiFPTG>(q, input, output,
ONEAPI::bit_and<int>(), ~0);
#endif // SPIRV_1_3

std::cout << "Test passed." << std::endl;
Expand Down
4 changes: 2 additions & 2 deletions SYCL/SubGroup/broadcast.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,8 +23,8 @@ template <typename T> void check(queue &Queue) {
ONEAPI::sub_group SG = NdItem.get_sub_group();
/*Broadcast GID of element with SGLID == SGID % SGMLR*/
syclacc[NdItem.get_global_id()] =
broadcast(SG, T(NdItem.get_global_id(0)),
SG.get_group_id() % SG.get_max_local_range()[0]);
ONEAPI::broadcast(SG, T(NdItem.get_global_id(0)),
SG.get_group_id() % SG.get_max_local_range()[0]);
if (NdItem.get_global_id(0) == 0)
sgsizeacc[0] = SG.get_max_local_range()[0];
});
Expand Down
4 changes: 2 additions & 2 deletions SYCL/SubGroup/generic_reduce.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,10 +25,10 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
ONEAPI::sub_group sg = NdItem.get_sub_group();
if (skip_init) {
acc[NdItem.get_global_id(0)] =
reduce(sg, T(NdItem.get_global_id(0)), op);
ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), op);
} else {
acc[NdItem.get_global_id(0)] =
reduce(sg, T(NdItem.get_global_id(0)), init, op);
ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), init, op);
}
if (NdItem.get_global_id(0) == 0)
sgsizeacc[0] = sg.get_max_local_range()[0];
Expand Down
4 changes: 2 additions & 2 deletions SYCL/SubGroup/reduce.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,10 +28,10 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
ONEAPI::sub_group sg = NdItem.get_sub_group();
if (skip_init) {
acc[NdItem.get_global_id(0)] =
reduce(sg, T(NdItem.get_global_id(0)), op);
ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), op);
} else {
acc[NdItem.get_global_id(0)] =
reduce(sg, T(NdItem.get_global_id(0)), init, op);
ONEAPI::reduce(sg, T(NdItem.get_global_id(0)), init, op);
}
if (NdItem.get_global_id(0) == 0)
sgsizeacc[0] = sg.get_max_local_range()[0];
Expand Down
12 changes: 6 additions & 6 deletions SYCL/SubGroup/scan.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,14 +30,14 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
ONEAPI::sub_group sg = NdItem.get_sub_group();
if (skip_init) {
exacc[NdItem.get_global_id(0)] =
exclusive_scan(sg, T(NdItem.get_global_id(0)), op);
ONEAPI::exclusive_scan(sg, T(NdItem.get_global_id(0)), op);
inacc[NdItem.get_global_id(0)] =
inclusive_scan(sg, T(NdItem.get_global_id(0)), op);
ONEAPI::inclusive_scan(sg, T(NdItem.get_global_id(0)), op);
} else {
exacc[NdItem.get_global_id(0)] =
exclusive_scan(sg, T(NdItem.get_global_id(0)), init, op);
inacc[NdItem.get_global_id(0)] =
inclusive_scan(sg, T(NdItem.get_global_id(0)), op, init);
exacc[NdItem.get_global_id(0)] = ONEAPI::exclusive_scan(
sg, T(NdItem.get_global_id(0)), init, op);
inacc[NdItem.get_global_id(0)] = ONEAPI::inclusive_scan(
sg, T(NdItem.get_global_id(0)), op, init);
}
if (NdItem.get_global_id(0) == 0)
sgsizeacc[0] = sg.get_max_local_range()[0];
Expand Down
4 changes: 2 additions & 2 deletions SYCL/SubGroup/vote.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,12 +48,12 @@ void check(queue Queue, const int G, const int L, const int D, const int R) {
cgh.parallel_for<class subgr>(NdRange, [=](nd_item<1> NdItem) {
ONEAPI::sub_group SG = NdItem.get_sub_group();
/* Set to 1 if any local ID in subgroup devided by D has remainder R */
if (any_of(SG, SG.get_local_id().get(0) % D == R)) {
if (ONEAPI::any_of(SG, SG.get_local_id().get(0) % D == R)) {
sganyacc[NdItem.get_global_id()] = 1;
}
/* Set to 1 if remainder of division of subgroup local ID by D is less
* than R for all work items in subgroup */
if (all_of(SG, SG.get_local_id().get(0) % D < R)) {
if (ONEAPI::all_of(SG, SG.get_local_id().get(0) % D < R)) {
sgallacc[NdItem.get_global_id()] = 1;
}
});
Expand Down