From 2c41b1d050db3f64d1dc50fa330c0ccfdddec7f8 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 21 Feb 2023 21:07:15 -0800 Subject: [PATCH 01/11] [SYCL] Test reduction for complex numbers. --- SYCL/Reduction/reduction_complex_nums.cpp | 52 +++++++++++++++++++++++ 1 file changed, 52 insertions(+) create mode 100644 SYCL/Reduction/reduction_complex_nums.cpp diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp new file mode 100644 index 0000000000..5871cce8e7 --- /dev/null +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -0,0 +1,52 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +#include +#include + +#include + +using namespace sycl; + +// Currently, Identityless reduction for complex numbers is +// only valid for plus operator. +// TODO: Extend this test case once we support known_identity for std::complex +// and more operators (apart from plus). +template +void test_identityless_reduction_for_complex_nums(queue &q) { + // Allocate and initialize buffer on the host with all 1's. + buffer> valuesBuf { 1024 }; + { + host_accessor a { valuesBuf }; + std::fill(a.begin(), a.end(), std::complex(1)); + } + + // Buffer to hold the reduction results. + std::complex sumResult = 0; + buffer> sumBuf { &sumResult, 1 }; + + q.submit([&](handler& cgh) { + auto inputVals = valuesBuf.template get_access(cgh); + auto sumReduction = reduction(sumBuf, cgh, plus>()); + + cgh.parallel_for(range<1> {1024}, sumReduction, + [=](id<1> idx, auto& sum) { + sum += inputVals[idx]; + }); + }); + + assert(sumBuf.get_host_access()[0] == static_cast(1024)); +} + +int main() { + queue q; + + test_identityless_reduction_for_complex_nums(q); + test_identityless_reduction_for_complex_nums(q); + test_identityless_reduction_for_complex_nums(q); + test_identityless_reduction_for_complex_nums(q); + + return 0; +} From 0d1a4777b30d69f075d81d0a21cda076a1a31a0c Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 22 Feb 2023 08:33:55 -0800 Subject: [PATCH 02/11] Fix formatting. --- SYCL/Reduction/reduction_complex_nums.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 5871cce8e7..3e67f47b9d 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -17,24 +17,22 @@ using namespace sycl; template void test_identityless_reduction_for_complex_nums(queue &q) { // Allocate and initialize buffer on the host with all 1's. - buffer> valuesBuf { 1024 }; + buffer> valuesBuf{1024}; { - host_accessor a { valuesBuf }; + host_accessor a{valuesBuf}; std::fill(a.begin(), a.end(), std::complex(1)); } // Buffer to hold the reduction results. std::complex sumResult = 0; - buffer> sumBuf { &sumResult, 1 }; + buffer> sumBuf{&sumResult, 1}; - q.submit([&](handler& cgh) { + q.submit([&](handler &cgh) { auto inputVals = valuesBuf.template get_access(cgh); auto sumReduction = reduction(sumBuf, cgh, plus>()); - cgh.parallel_for(range<1> {1024}, sumReduction, - [=](id<1> idx, auto& sum) { - sum += inputVals[idx]; - }); + cgh.parallel_for(range<1>{1024}, sumReduction, + [=](id<1> idx, auto &sum) { sum += inputVals[idx]; }); }); assert(sumBuf.get_host_access()[0] == static_cast(1024)); From 1482780d10138effef18c2ce4af363ca5f26968a Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 22 Feb 2023 09:15:17 -0800 Subject: [PATCH 03/11] Update test case to only check for std::complex and std::complex. --- SYCL/Reduction/reduction_complex_nums.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 3e67f47b9d..cd0d5b24c2 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -41,8 +41,6 @@ void test_identityless_reduction_for_complex_nums(queue &q) { int main() { queue q; - test_identityless_reduction_for_complex_nums(q); - test_identityless_reduction_for_complex_nums(q); test_identityless_reduction_for_complex_nums(q); test_identityless_reduction_for_complex_nums(q); From 3c99d052eabee2d1f84e5491cc44e01702973f7c Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 22 Feb 2023 11:06:20 -0800 Subject: [PATCH 04/11] Minor changes. --- SYCL/Reduction/reduction_complex_nums.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index cd0d5b24c2..39bfa83abd 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -20,7 +20,9 @@ void test_identityless_reduction_for_complex_nums(queue &q) { buffer> valuesBuf{1024}; { host_accessor a{valuesBuf}; - std::fill(a.begin(), a.end(), std::complex(1)); + T n = 0; + std::generate(a.begin(), a.end(), + [&n] { return std::complex(n, ++n + 1); }); } // Buffer to hold the reduction results. @@ -28,14 +30,14 @@ void test_identityless_reduction_for_complex_nums(queue &q) { buffer> sumBuf{&sumResult, 1}; q.submit([&](handler &cgh) { - auto inputVals = valuesBuf.template get_access(cgh); + accessor inputVals{valuesBuf, cgh, sycl::read_only}; auto sumReduction = reduction(sumBuf, cgh, plus>()); cgh.parallel_for(range<1>{1024}, sumReduction, [=](id<1> idx, auto &sum) { sum += inputVals[idx]; }); }); - assert(sumBuf.get_host_access()[0] == static_cast(1024)); + assert(sumBuf.get_host_access()[0] == std::complex(523776, 525824)); } int main() { From 9d8f19f089b1a8b8dc12699aa8ca905627f59f81 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 7 Mar 2023 08:22:11 -0800 Subject: [PATCH 05/11] Fix test failure on windows. --- SYCL/Reduction/reduction_complex_nums.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 39bfa83abd..91dc7a64a8 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -5,6 +5,7 @@ #include #include +#include #include From 5d5dab8752aecccdf1b1799dad28a60bea52b01a Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 7 Mar 2023 08:36:37 -0800 Subject: [PATCH 06/11] Minor formatting changes. --- SYCL/Reduction/reduction_complex_nums.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 91dc7a64a8..561f1c3990 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -3,9 +3,9 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out +#include #include #include -#include #include From 06c2de80fc66afb8b599904ab2713a53a4f50cfb Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 14 Mar 2023 07:51:03 -0700 Subject: [PATCH 07/11] Replace range with nd_range --- SYCL/Reduction/reduction_complex_nums.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 561f1c3990..c3544275e6 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -18,7 +18,7 @@ using namespace sycl; template void test_identityless_reduction_for_complex_nums(queue &q) { // Allocate and initialize buffer on the host with all 1's. - buffer> valuesBuf{1024}; + buffer> valuesBuf{255}; { host_accessor a{valuesBuf}; T n = 0; @@ -34,18 +34,21 @@ void test_identityless_reduction_for_complex_nums(queue &q) { accessor inputVals{valuesBuf, cgh, sycl::read_only}; auto sumReduction = reduction(sumBuf, cgh, plus>()); - cgh.parallel_for(range<1>{1024}, sumReduction, - [=](id<1> idx, auto &sum) { sum += inputVals[idx]; }); + cgh.parallel_for(nd_range<1>{255, 255}, sumReduction, + [=](nd_item<1> idx, auto &sum) { + sum += inputVals[idx.get_global_id(0)]; + }); }); - assert(sumBuf.get_host_access()[0] == std::complex(523776, 525824)); + assert(sumBuf.get_host_access()[0] == std::complex(32385, 32895)); } int main() { queue q; test_identityless_reduction_for_complex_nums(q); - test_identityless_reduction_for_complex_nums(q); + if (q.get_device().has(aspect::fp64)) + test_identityless_reduction_for_complex_nums(q); return 0; } From ef48bdf340ae71b9ae3ca80ad97c2d06309eafe5 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 21 Mar 2023 17:11:25 -0700 Subject: [PATCH 08/11] Fix typo. --- SYCL/Reduction/reduction_complex_nums.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index c3544275e6..c93c5613de 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -11,6 +11,8 @@ using namespace sycl; +#define BUFFER_SIZE 255 + // Currently, Identityless reduction for complex numbers is // only valid for plus operator. // TODO: Extend this test case once we support known_identity for std::complex @@ -18,12 +20,12 @@ using namespace sycl; template void test_identityless_reduction_for_complex_nums(queue &q) { // Allocate and initialize buffer on the host with all 1's. - buffer> valuesBuf{255}; + buffer> valuesBuf{BUFFER_SIZE}; { host_accessor a{valuesBuf}; T n = 0; std::generate(a.begin(), a.end(), - [&n] { return std::complex(n, ++n + 1); }); + [&n] { n++; return std::complex(n, n + 1); }); } // Buffer to hold the reduction results. @@ -34,7 +36,7 @@ void test_identityless_reduction_for_complex_nums(queue &q) { accessor inputVals{valuesBuf, cgh, sycl::read_only}; auto sumReduction = reduction(sumBuf, cgh, plus>()); - cgh.parallel_for(nd_range<1>{255, 255}, sumReduction, + cgh.parallel_for(nd_range<1>{BUFFER_SIZE, BUFFER_SIZE}, sumReduction, [=](nd_item<1> idx, auto &sum) { sum += inputVals[idx.get_global_id(0)]; }); From f33c528fb83cc3cb1a7ca83cd75eb2eba3e6b575 Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 21 Mar 2023 17:14:59 -0700 Subject: [PATCH 09/11] Fix assert statement. --- SYCL/Reduction/reduction_complex_nums.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index c93c5613de..9b5ca152f4 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -42,7 +42,7 @@ void test_identityless_reduction_for_complex_nums(queue &q) { }); }); - assert(sumBuf.get_host_access()[0] == std::complex(32385, 32895)); + assert(sumBuf.get_host_access()[0] == std::complex(32640,32895)); } int main() { From 297a086ef8cc5cd0c1a20b734f6f7936e2ca127b Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Tue, 21 Mar 2023 20:01:46 -0700 Subject: [PATCH 10/11] Fix formatting. --- SYCL/Reduction/reduction_complex_nums.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 9b5ca152f4..94bbb8ac6c 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -24,8 +24,10 @@ void test_identityless_reduction_for_complex_nums(queue &q) { { host_accessor a{valuesBuf}; T n = 0; - std::generate(a.begin(), a.end(), - [&n] { n++; return std::complex(n, n + 1); }); + std::generate(a.begin(), a.end(), [&n] { + n++; + return std::complex(n, n + 1); + }); } // Buffer to hold the reduction results. From cc9056047b12dd32bd11ff43258009554e50ea4b Mon Sep 17 00:00:00 2001 From: "Agarwal, Udit" Date: Wed, 22 Mar 2023 07:58:08 -0700 Subject: [PATCH 11/11] More formatting changes. --- SYCL/Reduction/reduction_complex_nums.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/Reduction/reduction_complex_nums.cpp b/SYCL/Reduction/reduction_complex_nums.cpp index 94bbb8ac6c..ec9432d19d 100644 --- a/SYCL/Reduction/reduction_complex_nums.cpp +++ b/SYCL/Reduction/reduction_complex_nums.cpp @@ -44,7 +44,7 @@ void test_identityless_reduction_for_complex_nums(queue &q) { }); }); - assert(sumBuf.get_host_access()[0] == std::complex(32640,32895)); + assert(sumBuf.get_host_access()[0] == std::complex(32640, 32895)); } int main() {