From 2b56445aa3230938ff3884b151103dfc884b0101 Mon Sep 17 00:00:00 2001 From: Marcel Koch Date: Fri, 20 Jan 2023 12:48:36 +0100 Subject: [PATCH 1/3] adds test for reduction on zero size inputs --- test/base/kernel_launch_generic.cpp | 173 +++++++++++++++------------- 1 file changed, 95 insertions(+), 78 deletions(-) diff --git a/test/base/kernel_launch_generic.cpp b/test/base/kernel_launch_generic.cpp index a90a5ea6c70..4e57904a9d2 100644 --- a/test/base/kernel_launch_generic.cpp +++ b/test/base/kernel_launch_generic.cpp @@ -321,44 +321,45 @@ TEST_F(KernelLaunch, Runs2DDense) void run1d_reduction(std::shared_ptr exec) { - gko::array output{exec, 1}; + gko::array output{exec, {-1l}}; + auto run_reduction = [&](int64 init, size_type size) { + gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( + exec, + [] GKO_KERNEL(auto i, auto a, auto dummy) { + static_assert(is_same::value, "index"); + static_assert(is_same::value, "value"); + static_assert(is_same::value, "dummy"); + return i + 1; + }, + [] GKO_KERNEL(auto i, auto j) { return i + j; }, + [] GKO_KERNEL(auto j) { return j * 2; }, init, output.get_data(), + size, output, move_only_val); + }; - gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( - exec, - [] GKO_KERNEL(auto i, auto a, auto dummy) { - static_assert(is_same::value, "index"); - static_assert(is_same::value, "value"); - static_assert(is_same::value, "dummy"); - return i + 1; - }, - [] GKO_KERNEL(auto i, auto j) { return i + j; }, - [] GKO_KERNEL(auto j) { return j * 2; }, int64{}, output.get_data(), - size_type{100000}, output, move_only_val); + { + SCOPED_TRACE("Size 0"); + run_reduction(int64{1}, size_type{0}); - // 2 * sum i=0...99999 (i+1) - ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), 10000100000LL); + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{1}); + } - gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( - exec, - [] GKO_KERNEL(auto i, auto a, auto dummy) { - static_assert(is_same::value, "index"); - static_assert(is_same::value, "value"); - static_assert(is_same::value, "dummy"); - return i + 1; - }, - [] GKO_KERNEL(auto i, auto j) { - static_assert(is_same::value, "a"); - static_assert(is_same::value, "b"); - return i + j; - }, - [] GKO_KERNEL(auto j) { - static_assert(is_same::value, "value"); - return j * 2; - }, - int64{}, output.get_data(), size_type{100}, output, move_only_val); + { + SCOPED_TRACE("Size 100000"); + run_reduction(int64{0}, size_type{100000}); - // 2 * sum i=0...99 (i+1) - ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), 10100LL); + // 2 * sum i=0...99999 (i+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), + int64{10000100000}); + } + + { + SCOPED_TRACE("Size 100"); + run_reduction(int64{0}, size_type{100}); + + // 2 * sum i=0...99 (i+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), + int64{10100}); + } } TEST_F(KernelLaunch, Reduction1D) { run1d_reduction(exec); } @@ -366,54 +367,70 @@ TEST_F(KernelLaunch, Reduction1D) { run1d_reduction(exec); } void run2d_reduction(std::shared_ptr exec) { - gko::array output{exec, 1}; + gko::array output{exec, {-1l}}; + auto run_reduction = [&](int64 init, gko::dim<2> size) { + gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( + exec, + [] GKO_KERNEL(auto i, auto j, auto a, auto dummy) { + static_assert(is_same::value, "index"); + static_assert(is_same::value, "index"); + static_assert(is_same::value, "value"); + static_assert(is_same::value, "dummy"); + return (i + 1) * (j + 1); + }, + [] GKO_KERNEL(auto i, auto j) { + static_assert(is_same::value, "a"); + static_assert(is_same::value, "b"); + return i + j; + }, + [] GKO_KERNEL(auto j) { + static_assert(is_same::value, "value"); + return j * 4; + }, + init, output.get_data(), size, output, move_only_val); + }; - gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( - exec, - [] GKO_KERNEL(auto i, auto j, auto a, auto dummy) { - static_assert(is_same::value, "index"); - static_assert(is_same::value, "index"); - static_assert(is_same::value, "value"); - static_assert(is_same::value, "dummy"); - return (i + 1) * (j + 1); - }, - [] GKO_KERNEL(auto i, auto j) { - static_assert(is_same::value, "a"); - static_assert(is_same::value, "b"); - return i + j; - }, - [] GKO_KERNEL(auto j) { - static_assert(is_same::value, "value"); - return j * 4; - }, - int64{}, output.get_data(), gko::dim<2>{1000, 100}, output, - move_only_val); + { + SCOPED_TRACE("Dim 0x0"); + run_reduction(int64{0}, gko::dim<2>{0, 0}); - // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) - ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), 10110100000LL); + // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); + } - gko::kernels::EXEC_NAMESPACE::run_kernel_reduction( - exec, - [] GKO_KERNEL(auto i, auto j, auto a, auto dummy) { - static_assert(is_same::value, "index"); - static_assert(is_same::value, "index"); - static_assert(is_same::value, "value"); - static_assert(is_same::value, "dummy"); - return (i + 1) * (j + 1); - }, - [] GKO_KERNEL(auto i, auto j) { - static_assert(is_same::value, "a"); - static_assert(is_same::value, "b"); - return i + j; - }, - [] GKO_KERNEL(auto j) { - static_assert(is_same::value, "value"); - return j * 4; - }, - int64{}, output.get_data(), gko::dim<2>{10, 10}, output, move_only_val); + { + SCOPED_TRACE("Dim 0x10"); + run_reduction(int64{0}, gko::dim<2>{0, 10}); + + // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); + } + + { + SCOPED_TRACE("Dim 10x0"); + run_reduction(int64{0}, gko::dim<2>{10, 0}); + + // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); + } + + { + SCOPED_TRACE("Dim 1000x100"); + run_reduction(int64{0}, gko::dim<2>{1000, 100}); + + // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), + int64{10110100000}); + } - // 4 * sum i=0...9 sum j=0...9 of (i+1)*(j+1) - ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), 12100LL); + { + SCOPED_TRACE("Dim 10x10"); + run_reduction(int64{0}, gko::dim<2>{10, 10}); + + // 4 * sum i=0...9 sum j=0...9 of (i+1)*(j+1) + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), + int64{12100}); + } } TEST_F(KernelLaunch, Reduction2D) { run2d_reduction(exec); } From a71660c86128956866387d11ad674393b7922628 Mon Sep 17 00:00:00 2001 From: Marcel Koch Date: Thu, 6 Jul 2023 12:07:52 +0200 Subject: [PATCH 2/3] fix dpcpp reduction for size=0 inputs --- dpcpp/base/kernel_launch_reduction.dp.hpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/dpcpp/base/kernel_launch_reduction.dp.hpp b/dpcpp/base/kernel_launch_reduction.dp.hpp index 6cae0c72dcb..1cf7c1f774a 100644 --- a/dpcpp/base/kernel_launch_reduction.dp.hpp +++ b/dpcpp/base/kernel_launch_reduction.dp.hpp @@ -194,8 +194,8 @@ void run_kernel_reduction_impl(std::shared_ptr exec, } else { queue->submit([&](sycl::handler& cgh) { generic_kernel_reduction_1d( - cgh, static_cast(size), num_workgroups, fn, op, finalize, - identity, result, args...); + cgh, static_cast(size), 1, fn, op, finalize, identity, + result, args...); }); } } @@ -240,9 +240,9 @@ void run_kernel_reduction_impl(std::shared_ptr exec, }); } else { queue->submit([&](sycl::handler& cgh) { - generic_kernel_reduction_2d( - cgh, rows, cols, num_workgroups, fn, op, finalize, identity, - result, args...); + generic_kernel_reduction_2d(cgh, rows, cols, 1, fn, + op, finalize, identity, + result, args...); }); } } From 187c25afe98542a7559e328c3473c8250e9b139a Mon Sep 17 00:00:00 2001 From: Marcel Koch Date: Fri, 7 Jul 2023 13:32:22 +0200 Subject: [PATCH 3/3] review updates: - use correct identity value - remove incorrect comments Co-authored-by: Tobias Ribizel --- test/base/kernel_launch_generic.cpp | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/test/base/kernel_launch_generic.cpp b/test/base/kernel_launch_generic.cpp index 4e57904a9d2..3dd1570c5f8 100644 --- a/test/base/kernel_launch_generic.cpp +++ b/test/base/kernel_launch_generic.cpp @@ -338,16 +338,15 @@ void run1d_reduction(std::shared_ptr exec) { SCOPED_TRACE("Size 0"); - run_reduction(int64{1}, size_type{0}); + run_reduction(int64{0}, size_type{0}); - ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{1}); + ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); } { SCOPED_TRACE("Size 100000"); run_reduction(int64{0}, size_type{100000}); - // 2 * sum i=0...99999 (i+1) ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{10000100000}); } @@ -394,7 +393,6 @@ void run2d_reduction(std::shared_ptr exec) SCOPED_TRACE("Dim 0x0"); run_reduction(int64{0}, gko::dim<2>{0, 0}); - // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); } @@ -402,7 +400,6 @@ void run2d_reduction(std::shared_ptr exec) SCOPED_TRACE("Dim 0x10"); run_reduction(int64{0}, gko::dim<2>{0, 10}); - // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); } @@ -410,7 +407,6 @@ void run2d_reduction(std::shared_ptr exec) SCOPED_TRACE("Dim 10x0"); run_reduction(int64{0}, gko::dim<2>{10, 0}); - // 4 * sum i=0...999 sum j=0...99 of (i+1)*(j+1) ASSERT_EQ(exec->copy_val_to_host(output.get_const_data()), int64{0}); }