From ebc797d77a2c567270cac8ee0f0a588736ebbf72 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 9 Nov 2022 15:08:35 -0800 Subject: [PATCH 1/5] [SYCL] Add a test targeting internal reduction APIs That way we can bypass strategy selection and test any of them with the same type/operation. --- SYCL/Reduction/reduction_internal.cpp | 131 ++++++++++++++++++++++++++ 1 file changed, 131 insertions(+) create mode 100644 SYCL/Reduction/reduction_internal.cpp diff --git a/SYCL/Reduction/reduction_internal.cpp b/SYCL/Reduction/reduction_internal.cpp new file mode 100644 index 0000000000..63bf59b24a --- /dev/null +++ b/SYCL/Reduction/reduction_internal.cpp @@ -0,0 +1,131 @@ +// 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 +using namespace sycl; + +template auto get_global_range(range Range) { return Range; } +template auto get_global_range(nd_range NDRange) { + return NDRange.get_global_range(); +} + +template +auto get_global_id(item Item) { + return Item.get_id(); +} +template auto get_global_id(nd_item NDItem) { + return NDItem.get_global_id(); +} + +template auto get_global_id(id Id) { return Id; } + +template +static void test(RangeTy Range) { + queue q; + + // We can select strategy explicitly so no need to test all combinations of + // types/operations. + using T = int; + using BinOpTy = std::plus; + + T Init{19}; + + auto Red = [&]() { + if constexpr (UseUSM) + return malloc_device(1, q); + else + return buffer{1}; + }(); + auto GetRedAcc = [&](handler &cgh) { + if constexpr (UseUSM) + return Red; + else + return accessor{Red, cgh}; + }; + + q.submit([&](handler &cgh) { + auto RedAcc = GetRedAcc(cgh); + cgh.single_task([=]() { RedAcc[0] = Init; }); + }).wait(); + + q.submit([&](handler &cgh) { + auto RedSycl = [&]() { + if constexpr (UseUSM) + if constexpr (InitToIdentity) + return reduction(Red, BinOpTy{}, + property::reduction::initialize_to_identity{}); + else + return reduction(Red, BinOpTy{}); + else if constexpr (InitToIdentity) + return reduction(Red, cgh, BinOpTy{}, + property::reduction::initialize_to_identity{}); + else + return reduction(Red, cgh, BinOpTy{}); + }(); + detail::reduction_parallel_for( + cgh, Range, ext::oneapi::experimental::detail::empty_properties_t{}, + RedSycl, [=](auto Item, auto &Red) { Red.combine(T{1}); }); + }).wait(); + + auto *Result = malloc_shared(1, q); + q.submit([&](handler &cgh) { + auto RedAcc = GetRedAcc(cgh); + cgh.single_task([=]() { *Result = RedAcc[0]; }); + }).wait(); + + auto N = get_global_range(Range).size(); + int Expected = InitToIdentity ? N : Init + N; +#ifdef __PRETTY_FUNCTION__ + std::cout << __PRETTY_FUNCTION__ << ": " << *Result << ", expected " + << Expected << std::endl; +#endif + assert(*Result == Expected); + + if constexpr (UseUSM) + free(Red, q); + free(Result, q); +} + +template +void loop_impl(std::integer_sequence, F &&f) { + (f(std::integral_constant{}), ...); +} + +template void loop(F &&f) { + loop_impl(std::make_integer_sequence{}, std::forward(f)); +} + +template +void testAllStrategies(RangeTy Range) { + loop<(int)detail::reduction::strategy::multi>([&](auto Id) { + constexpr auto Strategy = + // Skip auto_select == 0. + detail::reduction::strategy{decltype(Id)::value + 1}; + test(Range); + }); +} + +int main() { + auto TestRange = [](auto Range) { + testAllStrategies(Range); + testAllStrategies(Range); + testAllStrategies(Range); + testAllStrategies(Range); + }; + + TestRange(range<1>{42}); + TestRange(range<2>{8, 8}); + TestRange(range<3>{7, 7, 5}); + TestRange(nd_range<1>{range<1>{7}, range<1>{7}}); + TestRange(nd_range<1>{range<1>{3 * 3}, range<1>{3}}); + + // TODO: Strategies historically adopted from sycl::range implementation only + // support 1-Dim case. + // + // TestRange(nd_range<2>{range<2>{7, 3}, range<2> {7, 3}}); + // TestRange(nd_range<2>{range<2>{14, 9}, range<2> {7, 3}}); + return 0; +} From 41585d4e4f3be8d1a6f0043b71865e3a1ea7925d Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 22 Nov 2022 10:42:01 -0800 Subject: [PATCH 2/5] Some updates for prints Hopefully would help investigate CI timeouts. Also, can be kept in the final version. --- SYCL/Reduction/reduction_internal.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/SYCL/Reduction/reduction_internal.cpp b/SYCL/Reduction/reduction_internal.cpp index 63bf59b24a..e6a658cdc6 100644 --- a/SYCL/Reduction/reduction_internal.cpp +++ b/SYCL/Reduction/reduction_internal.cpp @@ -78,10 +78,12 @@ static void test(RangeTy Range) { auto N = get_global_range(Range).size(); int Expected = InitToIdentity ? N : Init + N; -#ifdef __PRETTY_FUNCTION__ - std::cout << __PRETTY_FUNCTION__ << ": " << *Result << ", expected " - << Expected << std::endl; +#if defined(__PRETTY_FUNCTION__) + std::cout << __PRETTY_FUNCTION__; +#elif defined(__FUNCSIG__) + std::cout << __FUNCSIG__; #endif + std::cout << ": " << *Result << ", expected " << Expected << std::endl; assert(*Result == Expected); if constexpr (UseUSM) From 60c1b68e7189cda7fc713cf1b5b224730a16a7cf Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Tue, 22 Nov 2022 14:37:51 -0800 Subject: [PATCH 3/5] Add time output --- SYCL/Reduction/reduction_internal.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/SYCL/Reduction/reduction_internal.cpp b/SYCL/Reduction/reduction_internal.cpp index e6a658cdc6..2faa728c2a 100644 --- a/SYCL/Reduction/reduction_internal.cpp +++ b/SYCL/Reduction/reduction_internal.cpp @@ -4,6 +4,7 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out #include +#include using namespace sycl; template auto get_global_range(range Range) { return Range; } @@ -78,6 +79,10 @@ static void test(RangeTy Range) { auto N = get_global_range(Range).size(); int Expected = InitToIdentity ? N : Init + N; + + std::time_t end_time = + std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); + std::cout << std::ctime(&end_time); #if defined(__PRETTY_FUNCTION__) std::cout << __PRETTY_FUNCTION__; #elif defined(__FUNCSIG__) @@ -111,6 +116,9 @@ void testAllStrategies(RangeTy Range) { } int main() { + std::time_t end_time = + std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); + std::cout << std::ctime(&end_time); auto TestRange = [](auto Range) { testAllStrategies(Range); testAllStrategies(Range); From d4474611454093c757fb54b23380e2f4180c18f0 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 23 Nov 2022 11:58:37 -0800 Subject: [PATCH 4/5] Revert "Add time output" This reverts commit 60c1b68e7189cda7fc713cf1b5b224730a16a7cf. --- SYCL/Reduction/reduction_internal.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/SYCL/Reduction/reduction_internal.cpp b/SYCL/Reduction/reduction_internal.cpp index 2faa728c2a..e6a658cdc6 100644 --- a/SYCL/Reduction/reduction_internal.cpp +++ b/SYCL/Reduction/reduction_internal.cpp @@ -4,7 +4,6 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out #include -#include using namespace sycl; template auto get_global_range(range Range) { return Range; } @@ -79,10 +78,6 @@ static void test(RangeTy Range) { auto N = get_global_range(Range).size(); int Expected = InitToIdentity ? N : Init + N; - - std::time_t end_time = - std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); - std::cout << std::ctime(&end_time); #if defined(__PRETTY_FUNCTION__) std::cout << __PRETTY_FUNCTION__; #elif defined(__FUNCSIG__) @@ -116,9 +111,6 @@ void testAllStrategies(RangeTy Range) { } int main() { - std::time_t end_time = - std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); - std::cout << std::ctime(&end_time); auto TestRange = [](auto Range) { testAllStrategies(Range); testAllStrategies(Range); From b034bac4fac054bd5c44e2e56cdba964313468b6 Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Wed, 23 Nov 2022 12:23:40 -0800 Subject: [PATCH 5/5] Workaround timeouts on Windows --- SYCL/Reduction/reduction_internal.cpp | 58 +++++++++++++++++---------- 1 file changed, 36 insertions(+), 22 deletions(-) diff --git a/SYCL/Reduction/reduction_internal.cpp b/SYCL/Reduction/reduction_internal.cpp index e6a658cdc6..d87a6e8b62 100644 --- a/SYCL/Reduction/reduction_internal.cpp +++ b/SYCL/Reduction/reduction_internal.cpp @@ -21,24 +21,37 @@ template auto get_global_id(nd_item NDItem) { template auto get_global_id(id Id) { return Id; } +// We can select strategy explicitly so no need to test all combinations of +// types/operations. +using T = int; +using BinOpTy = std::plus; + +// On Windows, allocating new memory and then initializing it is slow for some +// reason (not related to reductions). Try to re-use the same memory between +// test cases. +struct RedStorage { + RedStorage(queue &q) : q(q), Ptr(malloc_device(1, q)), Buf(1) {} + ~RedStorage() { free(Ptr, q); } + + template auto get() { + if constexpr (UseUSM) + return Ptr; + else + return Buf; + } + queue &q; + T *Ptr; + buffer Buf; +}; + template -static void test(RangeTy Range) { - queue q; - - // We can select strategy explicitly so no need to test all combinations of - // types/operations. - using T = int; - using BinOpTy = std::plus; +static void test(RedStorage &Storage, RangeTy Range) { + queue &q = Storage.q; T Init{19}; - auto Red = [&]() { - if constexpr (UseUSM) - return malloc_device(1, q); - else - return buffer{1}; - }(); + auto Red = Storage.get(); auto GetRedAcc = [&](handler &cgh) { if constexpr (UseUSM) return Red; @@ -86,8 +99,6 @@ static void test(RangeTy Range) { std::cout << ": " << *Result << ", expected " << Expected << std::endl; assert(*Result == Expected); - if constexpr (UseUSM) - free(Red, q); free(Result, q); } @@ -101,21 +112,24 @@ template void loop(F &&f) { } template -void testAllStrategies(RangeTy Range) { +void testAllStrategies(RedStorage &Storage, RangeTy Range) { loop<(int)detail::reduction::strategy::multi>([&](auto Id) { constexpr auto Strategy = // Skip auto_select == 0. detail::reduction::strategy{decltype(Id)::value + 1}; - test(Range); + test(Storage, Range); }); } int main() { - auto TestRange = [](auto Range) { - testAllStrategies(Range); - testAllStrategies(Range); - testAllStrategies(Range); - testAllStrategies(Range); + queue q; + RedStorage Storage(q); + + auto TestRange = [&](auto Range) { + testAllStrategies(Storage, Range); + testAllStrategies(Storage, Range); + testAllStrategies(Storage, Range); + testAllStrategies(Storage, Range); }; TestRange(range<1>{42});