From 5b5ddac4a590b15501e19242c65eb5da377b5c4c Mon Sep 17 00:00:00 2001 From: Marcel Koch Date: Fri, 14 Feb 2025 10:56:11 +0100 Subject: [PATCH] [dist] review updates: - fix binary search usage in cuda/hip - refactor - add tests for segmented array assertion Co-authored-by: Tobias Ribizel Co-authored-by: Pratik Nayak Co-authored-by: Yu-Hsiang M. Tsai --- .../distributed/index_map_kernels.cpp | 37 ++++++----- core/distributed/index_map.cpp | 8 +-- core/distributed/index_map_kernels.hpp | 4 +- core/test/utils/assertions.hpp | 25 +++++-- core/test/utils/assertions_test.cpp | 48 +++++++++++++- dpcpp/distributed/index_map_kernels.dp.cpp | 4 +- dpcpp/distributed/partition_kernels.dp.cpp | 5 +- include/ginkgo/core/distributed/index_map.hpp | 7 +- omp/distributed/index_map_kernels.cpp | 44 +++++-------- reference/distributed/index_map_kernels.cpp | 30 ++++----- .../test/distributed/index_map_kernels.cpp | 58 +++++++++-------- test/distributed/index_map_kernels.cpp | 65 ------------------- 12 files changed, 156 insertions(+), 179 deletions(-) diff --git a/common/cuda_hip/distributed/index_map_kernels.cpp b/common/cuda_hip/distributed/index_map_kernels.cpp index fde1efbabb4..d5eb480ff92 100644 --- a/common/cuda_hip/distributed/index_map_kernels.cpp +++ b/common/cuda_hip/distributed/index_map_kernels.cpp @@ -302,20 +302,20 @@ void map_to_global( device_partition partition, device_segmented_array remote_global_idxs, experimental::distributed::comm_index_type rank, - const array& local_ids, + const array& local_idxs, experimental::distributed::index_space is, - array& global_ids) + array& global_idxs) { auto range_bounds = partition.offsets_begin; auto starting_indices = partition.starting_indices_begin; const auto& ranges_by_part = partition.ranges_by_part; - auto local_ids_it = local_ids.get_const_data(); - auto input_size = local_ids.get_size(); + auto local_idxs_it = local_idxs.get_const_data(); + auto input_size = local_idxs.get_size(); auto policy = thrust_policy(exec); - global_ids.resize_and_reset(local_ids.get_size()); - auto global_ids_it = global_ids.get_data(); + global_idxs.resize_and_reset(local_idxs.get_size()); + auto global_idxs_it = global_idxs.get_data(); auto map_local = [rank, ranges_by_part, range_bounds, starting_indices, partition] __device__(auto lid) { @@ -330,11 +330,16 @@ void map_to_global( auto local_ranges_size = static_cast(local_ranges.end - local_ranges.begin); - auto it = binary_search(int64(0), local_ranges_size, [=](const auto i) { - return starting_indices[local_ranges.begin[i]] >= lid; - }); + // the binary search finds the first local range, such that the starting + // index is larger than lid, thus lid is contained in the local range + // before that one auto local_range_id = - it != local_ranges_size ? it : max(int64(0), it - 1); + binary_search(int64(0), local_ranges_size, + [=](const auto i) { + return starting_indices[local_ranges.begin[i]] > + lid; + }) - + 1; auto range_id = local_ranges.begin[local_range_id]; return static_cast(lid - starting_indices[range_id]) + @@ -363,16 +368,16 @@ void map_to_global( }; if (is == experimental::distributed::index_space::local) { - thrust::transform(policy, local_ids_it, local_ids_it + input_size, - global_ids_it, map_local); + thrust::transform(policy, local_idxs_it, local_idxs_it + input_size, + global_idxs_it, map_local); } if (is == experimental::distributed::index_space::non_local) { - thrust::transform(policy, local_ids_it, local_ids_it + input_size, - global_ids_it, map_non_local); + thrust::transform(policy, local_idxs_it, local_idxs_it + input_size, + global_idxs_it, map_non_local); } if (is == experimental::distributed::index_space::combined) { - thrust::transform(policy, local_ids_it, local_ids_it + input_size, - global_ids_it, map_combined); + thrust::transform(policy, local_idxs_it, local_idxs_it + input_size, + global_idxs_it, map_combined); } } diff --git a/core/distributed/index_map.cpp b/core/distributed/index_map.cpp index e06c57c40a0..edd1a3cc3cb 100644 --- a/core/distributed/index_map.cpp +++ b/core/distributed/index_map.cpp @@ -93,15 +93,15 @@ array index_map::map_to_local( template array index_map::map_to_global( - const array& local_ids, index_space index_space_v) const + const array& local_idxs, index_space index_space_v) const { - array global_ids(exec_); + array global_idxs(exec_); exec_->run(index_map_kernels::make_map_to_global( to_device_const(partition_.get()), to_device(remote_global_idxs_), - rank_, local_ids, index_space_v, global_ids)); + rank_, local_idxs, index_space_v, global_idxs)); - return global_ids; + return global_idxs; } diff --git a/core/distributed/index_map_kernels.hpp b/core/distributed/index_map_kernels.hpp index c1d3be561f8..213810fd0ab 100644 --- a/core/distributed/index_map_kernels.hpp +++ b/core/distributed/index_map_kernels.hpp @@ -95,8 +95,8 @@ namespace kernels { device_partition partition, \ device_segmented_array remote_global_idxs, \ experimental::distributed::comm_index_type rank, \ - const array<_ltype>& local_ids, \ - experimental::distributed::index_space is, array<_gtype>& global_ids) + const array<_ltype>& local_idxs, \ + experimental::distributed::index_space is, array<_gtype>& global_idxs) #define GKO_DECLARE_ALL_AS_TEMPLATES \ diff --git a/core/test/utils/assertions.hpp b/core/test/utils/assertions.hpp index a74cae9daab..7f37ae108b8 100644 --- a/core/test/utils/assertions.hpp +++ b/core/test/utils/assertions.hpp @@ -24,6 +24,7 @@ #include #include #include +#include #include #include "core/base/batch_utilities.hpp" @@ -1014,12 +1015,6 @@ ::testing::AssertionResult segmented_array_equal( second.get_const_flat_data()) .copy_to_array(); - auto buffer_result = array_equal(first_expression, second_expression, - view_first, view_second); - if (buffer_result == ::testing::AssertionFailure()) { - return buffer_result << "Buffers of the segmented arrays mismatch"; - } - auto offsets_result = array_equal(first_expression, second_expression, first.get_offsets(), second.get_offsets()); @@ -1027,6 +1022,12 @@ ::testing::AssertionResult segmented_array_equal( return offsets_result << "Offsets of the segmented arrays mismatch"; } + auto buffer_result = array_equal(first_expression, second_expression, + view_first, view_second); + if (buffer_result == ::testing::AssertionFailure()) { + return buffer_result << "Buffers of the segmented arrays mismatch"; + } + return ::testing::AssertionSuccess(); } @@ -1414,6 +1415,18 @@ T* plain_ptr(T* ptr) } +/** + * Checks if two `gko::segmented_array`s are equal. + * + * Both the flat array buffer and the offsets of both arrays are tested + * for equality. + * + * Has to be called from within a google test unit test. + * Internally calls gko::test::assertions::segmented_array_equal(). + * + * @param _array1 first segmented array + * @param _array2 second segmented array + */ #define GKO_ASSERT_SEGMENTED_ARRAY_EQ(_array1, _array2) \ { \ ASSERT_PRED_FORMAT2(::gko::test::assertions::segmented_array_equal, \ diff --git a/core/test/utils/assertions_test.cpp b/core/test/utils/assertions_test.cpp index 73900397fbe..3be80230932 100644 --- a/core/test/utils/assertions_test.cpp +++ b/core/test/utils/assertions_test.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors +// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors // // SPDX-License-Identifier: BSD-3-Clause @@ -8,6 +8,7 @@ #include +#include #include #include @@ -218,4 +219,49 @@ TEST_F(ArraysNear, CanUseShortNotation) } +class SegmentedArraysEqual : public ::testing::Test { +protected: + using array = gko::array; + using iarray = gko::array; + using segmented_array = gko::segmented_array; + + std::shared_ptr exec = gko::ReferenceExecutor::create(); + + segmented_array arr1 = segmented_array::create_from_sizes( + array{exec, {1, 2, 3, 4, 5}}, iarray{exec, {2, 1, 2}}); + segmented_array arr2 = segmented_array::create_from_sizes( + array{exec, {1, 2, 3, 4, 5}}, iarray{exec, {2, 1, 2}}); + segmented_array arr3 = segmented_array::create_from_sizes( + array{exec, {1, 2, 3, 5, 6}}, iarray{exec, {2, 1, 2}}); + segmented_array arr4 = segmented_array::create_from_sizes( + array{exec, {1, 2, 3, 4, 5}}, iarray{exec, {3, 2}}); + segmented_array arr5 = segmented_array::create_from_sizes( + array{exec, {1, 2, 3, 4, 5}}, iarray{exec, {1, 2, 2}}); +}; + + +TEST_F(SegmentedArraysEqual, SucceedsIfEqual) +{ + GKO_ASSERT_SEGMENTED_ARRAY_EQ(arr1, arr2); +} + + +TEST_F(SegmentedArraysEqual, FailsIfValuesDifferent) +{ + GKO_ASSERT_SEGMENTED_ARRAY_EQ(arr1, arr3); +} + + +TEST_F(SegmentedArraysEqual, FailsIfOffsetsDifferent1) +{ + GKO_ASSERT_SEGMENTED_ARRAY_EQ(arr1, arr4); +} + + +TEST_F(SegmentedArraysEqual, FailsIfOffsetsDifferent2) +{ + GKO_ASSERT_SEGMENTED_ARRAY_EQ(arr1, arr5); +} + + } // namespace diff --git a/dpcpp/distributed/index_map_kernels.dp.cpp b/dpcpp/distributed/index_map_kernels.dp.cpp index 65fce495b6a..b5b022add72 100644 --- a/dpcpp/distributed/index_map_kernels.dp.cpp +++ b/dpcpp/distributed/index_map_kernels.dp.cpp @@ -50,9 +50,9 @@ void map_to_global( device_partition partition, device_segmented_array remote_global_idxs, experimental::distributed::comm_index_type rank, - const array& local_ids, + const array& local_idxs, experimental::distributed::index_space is, - array& global_ids) GKO_NOT_IMPLEMENTED; + array& global_idxs) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_LOCAL_GLOBAL_INDEX_TYPE( GKO_DECLARE_INDEX_MAP_MAP_TO_GLOBAL); diff --git a/dpcpp/distributed/partition_kernels.dp.cpp b/dpcpp/distributed/partition_kernels.dp.cpp index 86968b81dc4..bbff46af3af 100644 --- a/dpcpp/distributed/partition_kernels.dp.cpp +++ b/dpcpp/distributed/partition_kernels.dp.cpp @@ -140,10 +140,7 @@ void build_ranges_by_part(std::shared_ptr exec, range_ids.resize_and_reset(num_ranges); auto range_ids_ptr = range_ids.get_data(); - // fill range_ids with 0,...,num_ranges - 1 - run_kernel( - exec, [] GKO_KERNEL(auto i, auto rid) { rid[i] = i; }, num_ranges, - range_ids_ptr); + components::fill_seq_array(exec, range_ids_ptr, num_ranges); oneapi::dpl::stable_sort(policy, range_ids_ptr, range_ids_ptr + num_ranges, [range_parts](const auto rid_a, const auto rid_b) { diff --git a/include/ginkgo/core/distributed/index_map.hpp b/include/ginkgo/core/distributed/index_map.hpp index 0c690eb462e..c8c402b9038 100644 --- a/include/ginkgo/core/distributed/index_map.hpp +++ b/include/ginkgo/core/distributed/index_map.hpp @@ -85,15 +85,16 @@ struct index_map { /** * Maps local indices to global indices * - * @param local_ids the local indices to map + * @param local_idxs the local indices to map * @param index_space_v the index space in which the passed-in local * indices are defined * * @return the mapped global indices. Any local index, that is not in the * specified index space is mapped to invalid_index */ - array map_to_global(const array& local_ids, - index_space index_space_v) const; + array map_to_global( + const array& local_idxs, + index_space index_space_v) const; /** * \brief get size of index_space::local diff --git a/omp/distributed/index_map_kernels.cpp b/omp/distributed/index_map_kernels.cpp index 135023fd7c0..2a1f113eecf 100644 --- a/omp/distributed/index_map_kernels.cpp +++ b/omp/distributed/index_map_kernels.cpp @@ -245,62 +245,50 @@ void map_to_global( device_partition partition, device_segmented_array remote_global_idxs, experimental::distributed::comm_index_type rank, - const array& local_ids, + const array& local_idxs, experimental::distributed::index_space is, - array& global_ids) + array& global_idxs) { const auto& ranges_by_part = partition.ranges_by_part; auto local_ranges = ranges_by_part.get_segment(rank); - global_ids.resize_and_reset(local_ids.get_size()); + global_idxs.resize_and_reset(local_idxs.get_size()); auto local_size = static_cast(partition.part_sizes_begin[rank]); auto remote_size = static_cast( remote_global_idxs.flat_end - remote_global_idxs.flat_begin); size_type local_range_id = 0; - if (is == experimental::distributed::index_space::local) { #pragma omp parallel for firstprivate(local_range_id) - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; + for (size_type i = 0; i < local_idxs.get_size(); ++i) { + auto lid = local_idxs.get_const_data()[i]; + if (is == experimental::distributed::index_space::local) { if (0 <= lid && lid < local_size) { local_range_id = find_local_range(lid, rank, partition, local_range_id); - global_ids.get_data()[i] = map_to_global( + global_idxs.get_data()[i] = map_to_global( lid, partition, local_ranges.begin[local_range_id]); } else { - global_ids.get_data()[i] = invalid_index(); + global_idxs.get_data()[i] = invalid_index(); } - } - } - if (is == experimental::distributed::index_space::non_local) { -#pragma omp parallel for - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; - + } else if (is == experimental::distributed::index_space::non_local) { if (0 <= lid && lid < remote_size) { - global_ids.get_data()[i] = remote_global_idxs.flat_begin[lid]; + global_idxs.get_data()[i] = remote_global_idxs.flat_begin[lid]; } else { - global_ids.get_data()[i] = invalid_index(); + global_idxs.get_data()[i] = invalid_index(); } - } - } - if (is == experimental::distributed::index_space::combined) { -#pragma omp parallel for firstprivate(local_range_id) - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; - + } else if (is == experimental::distributed::index_space::combined) { if (0 <= lid && lid < local_size) { local_range_id = find_local_range(lid, rank, partition, local_range_id); - global_ids.get_data()[i] = map_to_global( + global_idxs.get_data()[i] = map_to_global( lid, partition, local_ranges.begin[local_range_id]); } else if (local_size <= lid && lid < local_size + remote_size) { - global_ids.get_data()[i] = + global_idxs.get_data()[i] = remote_global_idxs.flat_begin[lid - local_size]; } else { - global_ids.get_data()[i] = invalid_index(); + global_idxs.get_data()[i] = invalid_index(); } } } @@ -308,8 +296,6 @@ void map_to_global( GKO_INSTANTIATE_FOR_EACH_LOCAL_GLOBAL_INDEX_TYPE( GKO_DECLARE_INDEX_MAP_MAP_TO_GLOBAL); - - } // namespace index_map } // namespace omp } // namespace kernels diff --git a/reference/distributed/index_map_kernels.cpp b/reference/distributed/index_map_kernels.cpp index ed7e9bdcc52..23a81dcc903 100644 --- a/reference/distributed/index_map_kernels.cpp +++ b/reference/distributed/index_map_kernels.cpp @@ -205,14 +205,14 @@ void map_to_global( device_partition partition, device_segmented_array remote_global_idxs, experimental::distributed::comm_index_type rank, - const array& local_ids, + const array& local_idxs, experimental::distributed::index_space is, - array& global_ids) + array& global_idxs) { const auto& ranges_by_part = partition.ranges_by_part; auto local_ranges = ranges_by_part.get_segment(rank); - global_ids.resize_and_reset(local_ids.get_size()); + global_idxs.resize_and_reset(local_idxs.get_size()); auto local_size = static_cast(partition.part_sizes_begin[rank]); @@ -246,22 +246,14 @@ void map_to_global( } }; - if (is == experimental::distributed::index_space::local) { - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; - global_ids.get_data()[i] = map_local(lid); - } - } - if (is == experimental::distributed::index_space::non_local) { - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; - global_ids.get_data()[i] = map_non_local(lid); - } - } - if (is == experimental::distributed::index_space::combined) { - for (size_type i = 0; i < local_ids.get_size(); ++i) { - auto lid = local_ids.get_const_data()[i]; - global_ids.get_data()[i] = map_combined(lid); + for (size_type i = 0; i < local_idxs.get_size(); ++i) { + auto lid = local_idxs.get_const_data()[i]; + if (is == experimental::distributed::index_space::local) { + global_idxs.get_data()[i] = map_local(lid); + } else if (is == experimental::distributed::index_space::non_local) { + global_idxs.get_data()[i] = map_non_local(lid); + } else if (is == experimental::distributed::index_space::combined) { + global_idxs.get_data()[i] = map_combined(lid); } } } diff --git a/reference/test/distributed/index_map_kernels.cpp b/reference/test/distributed/index_map_kernels.cpp index 525bf90b5de..504e013e452 100644 --- a/reference/test/distributed/index_map_kernels.cpp +++ b/reference/test/distributed/index_map_kernels.cpp @@ -200,68 +200,70 @@ TEST_F(IndexMap, CanGetLocalWithCombinedISWithInvalid) TEST_F(IndexMap, CanGetGlobalWithLocalISWithInvalid) { - gko::array global_ids(ref); - gko::array local_ids(ref, {5, 4, 10, 3, 2, 1, 0, 100, 4}); + gko::array global_idxs(ref); + gko::array local_idxs(ref, + {5, 4, 10, 3, 2, 1, 0, 100, 4}); auto remote_global_idxs = gko::segmented_array{ref}; gko::kernels::reference::index_map::map_to_global( ref, to_device_const(part_large.get()), - to_device_const(remote_global_idxs), 0, local_ids, - gko::experimental::distributed::index_space::local, global_ids); + to_device_const(remote_global_idxs), 0, local_idxs, + gko::experimental::distributed::index_space::local, global_idxs); auto invalid = gko::invalid_index(); gko::array expected( ref, I{17, 16, invalid, 15, 2, 1, 0, invalid, 16}); - GKO_ASSERT_ARRAY_EQ(global_ids, expected); + GKO_ASSERT_ARRAY_EQ(global_idxs, expected); } TEST_F(IndexMap, CanGetGlobalWithNonLocalISWithInvalid) { - gko::array global_ids(ref); - gko::array local_ids(ref, {5, 4, 10, 3, 2, 1, 0, 100, 4}); + gko::array global_idxs(ref); + gko::array local_idxs(ref, + {5, 4, 10, 3, 2, 1, 0, 100, 4}); auto remote_global_idxs = gko::segmented_array::create_from_sizes( {ref, {0, 1, 9, 8, 7, 6}}, {ref, {2, 4}}); gko::kernels::reference::index_map::map_to_global( ref, to_device_const(part_large.get()), - to_device_const(remote_global_idxs), 1, local_ids, - gko::experimental::distributed::index_space::non_local, global_ids); + to_device_const(remote_global_idxs), 1, local_idxs, + gko::experimental::distributed::index_space::non_local, global_idxs); auto invalid = gko::invalid_index(); gko::array expected( ref, I{6, 7, invalid, 8, 9, 1, 0, invalid, 7}); - GKO_ASSERT_ARRAY_EQ(global_ids, expected); + GKO_ASSERT_ARRAY_EQ(global_idxs, expected); } TEST_F(IndexMap, CanGetGlobalWithCombinedISWithInvalid) { - gko::array global_ids(ref); - gko::array local_ids(ref, {2, 5, 133, 6, 10}); + gko::array global_idxs(ref); + gko::array local_idxs(ref, {2, 5, 133, 6, 10}); auto remote_global_idxs = gko::segmented_array::create_from_sizes( {ref, {0, 1, 9, 8, 7, 6}}, {ref, {2, 4}}); gko::kernels::reference::index_map::map_to_global( ref, to_device_const(part_large.get()), - to_device_const(remote_global_idxs), 1, local_ids, - gko::experimental::distributed::index_space::combined, global_ids); + to_device_const(remote_global_idxs), 1, local_idxs, + gko::experimental::distributed::index_space::combined, global_idxs); auto invalid = gko::invalid_index(); gko::array expected( ref, I{5, 14, invalid, 0, 7}); - GKO_ASSERT_ARRAY_EQ(global_ids, expected); + GKO_ASSERT_ARRAY_EQ(global_idxs, expected); } TEST_F(IndexMap, RoundTripGlobalWithCombinedIS) { gko::array result(ref); - gko::array global_ids(ref, - {5, 14, 14, 0, 3, 4, 8, 7, 12}); - gko::array local_ids(ref); + gko::array global_idxs(ref, + {5, 14, 14, 0, 3, 4, 8, 7, 12}); + gko::array local_idxs(ref); auto remote_global_idxs = gko::segmented_array::create_from_sizes( {ref, {0, 1, 6, 7, 8, 9}}, {ref, {2, 4}}); @@ -269,22 +271,22 @@ TEST_F(IndexMap, RoundTripGlobalWithCombinedIS) gko::kernels::reference::index_map::map_to_local( ref, part_large.get(), remote_target_ids, - to_device_const(remote_global_idxs), 1, global_ids, - gko::experimental::distributed::index_space::combined, local_ids); + to_device_const(remote_global_idxs), 1, global_idxs, + gko::experimental::distributed::index_space::combined, local_idxs); gko::kernels::reference::index_map::map_to_global( ref, to_device_const(part_large.get()), - to_device_const(remote_global_idxs), 1, local_ids, + to_device_const(remote_global_idxs), 1, local_idxs, gko::experimental::distributed::index_space::combined, result); - GKO_ASSERT_ARRAY_EQ(result, global_ids); + GKO_ASSERT_ARRAY_EQ(result, global_idxs); } TEST_F(IndexMap, RoundTripLocalWithCombinedIS) { gko::array result(ref); - gko::array global_ids(ref); - gko::array local_ids(ref, {2, 5, 4, 6, 3, 3, 10}); + gko::array global_idxs(ref); + gko::array local_idxs(ref, {2, 5, 4, 6, 3, 3, 10}); auto remote_global_idxs = gko::segmented_array::create_from_sizes( {ref, {0, 1, 6, 7, 8, 9}}, {ref, {2, 4}}); @@ -292,12 +294,12 @@ TEST_F(IndexMap, RoundTripLocalWithCombinedIS) gko::kernels::reference::index_map::map_to_global( ref, to_device_const(part_large.get()), - to_device_const(remote_global_idxs), 1, local_ids, - gko::experimental::distributed::index_space::combined, global_ids); + to_device_const(remote_global_idxs), 1, local_idxs, + gko::experimental::distributed::index_space::combined, global_idxs); gko::kernels::reference::index_map::map_to_local( ref, part_large.get(), remote_target_ids, - to_device_const(remote_global_idxs), 1, global_ids, + to_device_const(remote_global_idxs), 1, global_idxs, gko::experimental::distributed::index_space::combined, result); - GKO_ASSERT_ARRAY_EQ(result, local_ids); + GKO_ASSERT_ARRAY_EQ(result, local_idxs); } diff --git a/test/distributed/index_map_kernels.cpp b/test/distributed/index_map_kernels.cpp index 1b266fbafd0..e1d8a56e98c 100644 --- a/test/distributed/index_map_kernels.cpp +++ b/test/distributed/index_map_kernels.cpp @@ -117,9 +117,6 @@ class IndexMap : public CommonTestFixture { IndexMap() { - std::random_device rd; - std::uniform_int_distribution<> engine_dist( - 0, std::numeric_limits::max()); engine.seed(490729788); auto connections = @@ -407,26 +404,6 @@ TEST_F(IndexMap, GetLocalWithCombinedIndexSpaceWithInvalidIndexSameAsRef) TEST_F(IndexMap, GetGlobalWithLocalIndexSpaceSameAsRef) -{ - auto query = generate_to_global_query(ref, local_size, 33); - auto dquery = gko::array(exec, query); - auto result = gko::array(ref); - auto dresult = gko::array(exec); - - gko::kernels::reference::index_map::map_to_global( - ref, to_device_const(part.get()), to_device_const(remote_global_idxs), - this_rank, query, gko::experimental::distributed::index_space::local, - result); - gko::kernels::GKO_DEVICE_NAMESPACE::index_map::map_to_global( - exec, to_device_const(dpart.get()), - to_device_const(dremote_global_idxs), this_rank, dquery, - gko::experimental::distributed::index_space::local, dresult); - - GKO_ASSERT_ARRAY_EQ(result, dresult); -} - - -TEST_F(IndexMap, GetGlobalWithLocalIndexSpaceWithInvalidIndexSameAsRef) { auto query = generate_to_global_query(ref, local_size * 2, 33); auto dquery = gko::array(exec, query); @@ -447,27 +424,6 @@ TEST_F(IndexMap, GetGlobalWithLocalIndexSpaceWithInvalidIndexSameAsRef) TEST_F(IndexMap, GetGlobalWithNonLocalIndexSpaceSameAsRef) -{ - auto query = - generate_to_global_query(ref, remote_global_idxs.get_size(), 33); - auto dquery = gko::array(exec, query); - auto result = gko::array(ref); - auto dresult = gko::array(exec); - - gko::kernels::reference::index_map::map_to_global( - ref, to_device_const(part.get()), to_device_const(remote_global_idxs), - this_rank, query, - gko::experimental::distributed::index_space::non_local, result); - gko::kernels::GKO_DEVICE_NAMESPACE::index_map::map_to_global( - exec, to_device_const(dpart.get()), - to_device_const(dremote_global_idxs), this_rank, dquery, - gko::experimental::distributed::index_space::non_local, dresult); - - GKO_ASSERT_ARRAY_EQ(result, dresult); -} - - -TEST_F(IndexMap, GetGlobalWithNonLocalIndexSpaceWithInvalidIndexSameAsRef) { auto query = generate_to_global_query(ref, remote_global_idxs.get_size() * 2, 33); @@ -489,27 +445,6 @@ TEST_F(IndexMap, GetGlobalWithNonLocalIndexSpaceWithInvalidIndexSameAsRef) TEST_F(IndexMap, GetGlobalWithCombinedIndexSpaceSameAsRef) -{ - auto query = generate_to_global_query( - ref, local_size + remote_global_idxs.get_size(), 33); - auto dquery = gko::array(exec, query); - auto result = gko::array(ref); - auto dresult = gko::array(exec); - - gko::kernels::reference::index_map::map_to_global( - ref, to_device_const(part.get()), to_device_const(remote_global_idxs), - this_rank, query, gko::experimental::distributed::index_space::combined, - result); - gko::kernels::GKO_DEVICE_NAMESPACE::index_map::map_to_global( - exec, to_device_const(dpart.get()), - to_device_const(dremote_global_idxs), this_rank, dquery, - gko::experimental::distributed::index_space::combined, dresult); - - GKO_ASSERT_ARRAY_EQ(result, dresult); -} - - -TEST_F(IndexMap, GetGlobalWithCombinedIndexSpaceWithInvalidIndexSameAsRef) { auto query = generate_to_global_query( ref, (local_size + remote_global_idxs.get_size()) * 2, 33);