diff --git a/common/cuda_hip/distributed/partition_helpers_kernels.hpp.inc b/common/cuda_hip/distributed/partition_helpers_kernels.hpp.inc new file mode 100644 index 00000000000..f92794ec138 --- /dev/null +++ b/common/cuda_hip/distributed/partition_helpers_kernels.hpp.inc @@ -0,0 +1,54 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +template +void sort_by_range_start( + std::shared_ptr exec, + array& range_start_ends, + array& part_ids) +{ + auto num_ranges = range_start_ends.get_num_elems() / 2; + auto strided_indices = thrust::make_transform_iterator( + thrust::make_counting_iterator(0), + [] __host__ __device__(const int i) { return 2 * i; }); + auto start_it = thrust::make_permutation_iterator( + range_start_ends.get_data(), strided_indices); + auto end_it = thrust::make_permutation_iterator( + range_start_ends.get_data() + 1, strided_indices); + auto zip_it = thrust::make_zip_iterator( + thrust::make_tuple(end_it, part_ids.get_data())); + thrust::stable_sort_by_key(thrust_policy(exec), start_it, + start_it + num_ranges, zip_it); +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START); diff --git a/common/unified/CMakeLists.txt b/common/unified/CMakeLists.txt index 5a37eb022f9..67fc839d6a7 100644 --- a/common/unified/CMakeLists.txt +++ b/common/unified/CMakeLists.txt @@ -6,6 +6,7 @@ set(UNIFIED_SOURCES components/format_conversion_kernels.cpp components/precision_conversion_kernels.cpp components/reduce_array_kernels.cpp + distributed/partition_helpers_kernels.cpp distributed/partition_kernels.cpp matrix/coo_kernels.cpp matrix/csr_kernels.cpp diff --git a/common/unified/distributed/partition_helpers_kernels.cpp b/common/unified/distributed/partition_helpers_kernels.cpp new file mode 100644 index 00000000000..3c041dd7e4b --- /dev/null +++ b/common/unified/distributed/partition_helpers_kernels.cpp @@ -0,0 +1,102 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/distributed/partition_helpers_kernels.hpp" + + +#include "common/unified/base/kernel_launch.hpp" +#include "common/unified/base/kernel_launch_reduction.hpp" + + +namespace gko { +namespace kernels { +namespace GKO_DEVICE_NAMESPACE { +namespace partition_helpers { + + +template +void check_consecutive_ranges(std::shared_ptr exec, + const array& range_start_ends, + bool& result) +{ + array result_uint32{exec, 1}; + auto num_ranges = range_start_ends.get_num_elems() / 2; + // need additional guard because DPCPP doesn't return the initial value for + // empty inputs + if (num_ranges > 1) { + run_kernel_reduction( + exec, + [] GKO_KERNEL(const auto i, const auto* ranges) { + return ranges[2 * i] == ranges[2 * i + 1]; + }, + [] GKO_KERNEL(const auto a, const auto b) { + return static_cast(a && b); + }, + [] GKO_KERNEL(auto x) { return x; }, static_cast(true), + result_uint32.get_data(), num_ranges - 1, + range_start_ends.get_const_data() + 1); + result = + static_cast(exec->copy_val_to_host(result_uint32.get_data())); + } else { + result = true; + } +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_CHECK_CONSECUTIVE_RANGES); + + +template +void compress_ranges(std::shared_ptr exec, + const array& range_start_ends, + array& range_offsets) +{ + run_kernel( + exec, + [] GKO_KERNEL(const auto i, const auto* start_ends, auto* offsets) { + if (i == 0) { + offsets[0] = start_ends[0]; + } + offsets[i + 1] = start_ends[2 * i + 1]; + }, + range_offsets.get_num_elems() - 1, range_start_ends.get_const_data(), + range_offsets.get_data()); +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_COMPRESS_RANGES); + + +} // namespace partition_helpers +} // namespace GKO_DEVICE_NAMESPACE +} // namespace kernels +} // namespace gko diff --git a/common/unified/distributed/partition_kernels.cpp b/common/unified/distributed/partition_kernels.cpp index cb0f4813da5..dc13fec9f1b 100644 --- a/common/unified/distributed/partition_kernels.cpp +++ b/common/unified/distributed/partition_kernels.cpp @@ -66,19 +66,22 @@ void count_ranges(std::shared_ptr exec, template void build_from_contiguous(std::shared_ptr exec, const array& ranges, + const array& part_id_mapping, GlobalIndexType* range_bounds, comm_index_type* part_ids) { run_kernel( exec, - [] GKO_KERNEL(auto i, auto ranges, auto bounds, auto ids) { + [] GKO_KERNEL(auto i, auto ranges, auto mapping, auto bounds, auto ids, + bool uses_mapping) { if (i == 0) { bounds[0] = 0; } bounds[i + 1] = ranges[i + 1]; - ids[i] = i; + ids[i] = uses_mapping ? mapping[i] : i; }, - ranges.get_num_elems() - 1, ranges, range_bounds, part_ids); + ranges.get_num_elems() - 1, ranges, part_id_mapping, range_bounds, + part_ids, part_id_mapping.get_num_elems() > 0); } GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_PARTITION_BUILD_FROM_CONTIGUOUS); diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 8ec4502d9c7..7932976d6c9 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -95,6 +95,7 @@ if(GINKGO_BUILD_MPI) PRIVATE mpi/exception.cpp distributed/matrix.cpp + distributed/partition_helpers.cpp distributed/vector.cpp distributed/preconditioner/schwarz.cpp) endif() diff --git a/core/base/copy_assignable.hpp b/core/base/copy_assignable.hpp new file mode 100644 index 00000000000..7f5e4125e10 --- /dev/null +++ b/core/base/copy_assignable.hpp @@ -0,0 +1,130 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#ifndef GKO_CORE_BASE_COPY_ASSIGNABLE_HPP_ +#define GKO_CORE_BASE_COPY_ASSIGNABLE_HPP_ + + +#include + + +namespace gko { +namespace detail { + + +template +class copy_assignable; + + +/** + * Helper class to make a type copy assignable. + * + * This class wraps an object of a type that has a copy constructor, but not + * a copy assignment. This is most often the case for lambdas. The wrapped + * object can then be copy assigned, by relying on the copy constructor. + * + * @tparam T type with a copy constructor + */ +template +class copy_assignable< + T, typename std::enable_if::value>::type> { +public: + copy_assignable() = default; + + copy_assignable(const copy_assignable& other) + { + if (this != &other) { + *this = other; + } + } + + copy_assignable(copy_assignable&& other) noexcept + { + if (this != &other) { + *this = std::move(other); + } + } + + copy_assignable(const T& obj) : obj_{new (buf)(T)(obj)} {} + + copy_assignable(T&& obj) : obj_{new (buf)(T)(std::move(obj))} {} + + copy_assignable& operator=(const copy_assignable& other) + { + if (this != &other) { + if (obj_) { + obj_->~T(); + } + obj_ = new (buf)(T)(*other.obj_); + } + return *this; + } + + copy_assignable& operator=(copy_assignable&& other) noexcept + { + if (this != &other) { + if (obj_) { + obj_->~T(); + } + obj_ = new (buf)(T)(std::move(*other.obj_)); + } + return *this; + } + + ~copy_assignable() + { + if (obj_) { + obj_->~T(); + } + } + + template + decltype(auto) operator()(Args&&... args) const + { + return (*obj_)(std::forward(args)...); + } + + T const& get() const { return *obj_; } + + T& get() { return *obj_; } + +private: + //!< Store wrapped object on the stack, should use std::optional in c++17 + T* obj_{}; + alignas(T) unsigned char buf[sizeof(T)]; +}; + + +} // namespace detail +} // namespace gko + +#endif // GKO_CORE_BASE_COPY_ASSIGNABLE_HPP_ diff --git a/core/base/iterator_factory.hpp b/core/base/iterator_factory.hpp index 7ebbc510f74..bbc1d3b4b2b 100644 --- a/core/base/iterator_factory.hpp +++ b/core/base/iterator_factory.hpp @@ -42,6 +42,9 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/base/copy_assignable.hpp" + + namespace gko { namespace detail { @@ -366,6 +369,133 @@ void swap(zip_iterator_reference a, } +/** + * Random access iterator that uses a function to transform the index. + * + * For a function `fn` and an underlying iterator `it`, accessing the + * permute_iterator at index `i` will result in accessing `it[fn(i)]`. + * + * @tparam IteratorType Underlying iterator, has to be random access. + * @tparam PermuteFn A function `difference_type -> difference_type` that + * transforms any given index. It doesn't have to be a strict + * permutation of indices (i.e. not bijective). + */ +template +class permute_iterator { +public: + using difference_type = std::ptrdiff_t; + using value_type = typename std::iterator_traits::value_type; + using pointer = typename std::iterator_traits::pointer; + using reference = typename std::iterator_traits::reference; + using iterator_category = std::random_access_iterator_tag; + + permute_iterator() = default; + + explicit permute_iterator(IteratorType it, PermuteFn perm) + : it_{std::move(it)}, idx_{}, perm_{std::move(perm)} + {} + + permute_iterator& operator+=(difference_type i) + { + idx_ += i; + return *this; + } + + permute_iterator& operator-=(difference_type i) { return *this += -i; } + + permute_iterator& operator++() { return *this += 1; } + + permute_iterator operator++(int) + { + auto tmp = *this; + ++(*this); + return tmp; + } + + permute_iterator& operator--() { return *this -= 1; } + + permute_iterator operator--(int) + { + auto tmp = *this; + --(*this); + return tmp; + } + + permute_iterator operator+(difference_type i) const + { + auto tmp = *this; + tmp += i; + return tmp; + } + + friend permute_iterator operator+(difference_type i, + const permute_iterator& iter) + { + return iter + i; + } + + permute_iterator operator-(difference_type i) const + { + auto tmp = *this; + tmp -= i; + return tmp; + } + + difference_type operator-(const permute_iterator& other) const + { + return idx_ - other.idx_; + } + + reference operator*() const { return it_[perm_(idx_)]; } + + reference operator[](difference_type i) const { return *(*this + i); } + + bool operator==(const permute_iterator& other) const + { + return idx_ == other.idx_; + } + + bool operator!=(const permute_iterator& other) const + { + return !(*this == other); + } + + bool operator<(const permute_iterator& other) const + { + return idx_ < other.idx_; + } + + bool operator<=(const permute_iterator& other) const + { + return idx_ <= other.idx_; + } + + bool operator>(const permute_iterator& other) const + { + return !(*this <= other); + } + + bool operator>=(const permute_iterator& other) const + { + return !(*this < other); + } + +private: + IteratorType it_; + difference_type idx_; + copy_assignable perm_; +}; + + +template +permute_iterator make_permute_iterator( + IteratorType it, PermutationFn perm) +{ + return permute_iterator{std::move(it), + std::move(perm)}; +} + + } // namespace detail } // namespace gko diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index 0f898b3ae73..c8bbd2e0a31 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -45,6 +45,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/components/prefix_sum_kernels.hpp" #include "core/components/reduce_array_kernels.hpp" #include "core/distributed/matrix_kernels.hpp" +#include "core/distributed/partition_helpers_kernels.hpp" #include "core/distributed/partition_kernels.hpp" #include "core/distributed/vector_kernels.hpp" #include "core/factorization/cholesky_kernels.hpp" @@ -255,6 +256,17 @@ GKO_STUB_LOCAL_GLOBAL_TYPE(GKO_DECLARE_PARTITION_IS_ORDERED); } // namespace partition +namespace partition_helpers { + + +GKO_STUB_INDEX_TYPE(GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START); +GKO_STUB_INDEX_TYPE(GKO_DECLARE_PARTITION_HELPERS_CHECK_CONSECUTIVE_RANGES); +GKO_STUB_INDEX_TYPE(GKO_DECLARE_PARTITION_HELPERS_COMPRESS_RANGES); + + +} // namespace partition_helpers + + namespace distributed_vector { diff --git a/core/distributed/partition.cpp b/core/distributed/partition.cpp index a1db99396e7..bfeb5e8c286 100644 --- a/core/distributed/partition.cpp +++ b/core/distributed/partition.cpp @@ -76,14 +76,21 @@ Partition::build_from_mapping( template std::unique_ptr> Partition::build_from_contiguous( - std::shared_ptr exec, const array& ranges) + std::shared_ptr exec, const array& ranges, + const array& part_ids) { + GKO_ASSERT(part_ids.get_num_elems() == 0 || + part_ids.get_num_elems() + 1 == ranges.get_num_elems()); + + array empty(exec); auto local_ranges = make_temporary_clone(exec, &ranges); + auto local_part_ids = make_temporary_clone( + exec, part_ids.get_num_elems() > 0 ? &part_ids : &empty); auto result = Partition::create( exec, static_cast(ranges.get_num_elems() - 1), ranges.get_num_elems() - 1); exec->run(partition::make_build_from_contiguous( - *local_ranges.get(), result->offsets_.get_data(), + *local_ranges, *local_part_ids, result->offsets_.get_data(), result->part_ids_.get_data())); result->finalize_construction(); return result; @@ -117,7 +124,7 @@ void Partition::finalize_construction() template -bool Partition::has_connected_parts() +bool Partition::has_connected_parts() const { return this->get_num_parts() - this->get_num_empty_parts() == this->get_num_ranges(); @@ -125,7 +132,7 @@ bool Partition::has_connected_parts() template -bool Partition::has_ordered_parts() +bool Partition::has_ordered_parts() const { if (this->has_connected_parts()) { auto exec = this->get_executor(); diff --git a/core/distributed/partition_helpers.cpp b/core/distributed/partition_helpers.cpp new file mode 100644 index 00000000000..b1fd1dd9bc5 --- /dev/null +++ b/core/distributed/partition_helpers.cpp @@ -0,0 +1,155 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include + + +#include + + +#include "core/components/fill_array_kernels.hpp" +#include "core/distributed/partition_helpers_kernels.hpp" + + +namespace gko { +namespace experimental { +namespace distributed { +namespace components { +namespace { + + +GKO_REGISTER_OPERATION(fill_seq_array, components::fill_seq_array); + + +} // namespace +} // namespace components + + +namespace partition_helpers { +namespace { + + +GKO_REGISTER_OPERATION(sort_by_range_start, + partition_helpers::sort_by_range_start); +GKO_REGISTER_OPERATION(check_consecutive_ranges, + partition_helpers::check_consecutive_ranges); +GKO_REGISTER_OPERATION(compress_ranges, partition_helpers::compress_ranges); + + +} // namespace +} // namespace partition_helpers + + +template +std::unique_ptr> +build_partition_from_local_range(std::shared_ptr exec, + mpi::communicator comm, span local_range) +{ + std::array range{ + static_cast(local_range.begin), + static_cast(local_range.end)}; + + // make all range_start_ends available on each rank + // note: not all combination of MPI + GPU library seem to support + // mixing host and device buffers, e.g. OpenMPI 4.0.5 and Rocm 4.0 + auto mpi_exec = exec->get_master(); + array ranges_start_end(mpi_exec, comm.size() * 2); + ranges_start_end.fill(invalid_index()); + comm.all_gather(mpi_exec, range.data(), 2, ranges_start_end.get_data(), 2); + ranges_start_end.set_executor(exec); + + // make_sort_by_range_start + array part_ids(exec, comm.size()); + exec->run(components::make_fill_seq_array(part_ids.get_data(), + part_ids.get_num_elems())); + exec->run(partition_helpers::make_sort_by_range_start(ranges_start_end, + part_ids)); + + // check for consistency + bool consecutive_ranges = false; + exec->run(partition_helpers::make_check_consecutive_ranges( + ranges_start_end, consecutive_ranges)); + if (!consecutive_ranges) { + GKO_INVALID_STATE("The partition contains gaps."); + } + + // join (now consecutive) starts and ends into combined array + array ranges(exec, comm.size() + 1); + exec->run( + partition_helpers::make_compress_ranges(ranges_start_end, ranges)); + + return Partition::build_from_contiguous( + exec, ranges, part_ids); +} + +#define GKO_DECLARE_BUILD_PARTITION_FROM_LOCAL_RANGE(_local_type, \ + _global_type) \ + std::unique_ptr> \ + build_partition_from_local_range(std::shared_ptr exec, \ + mpi::communicator comm, span local_range) +GKO_INSTANTIATE_FOR_EACH_LOCAL_GLOBAL_INDEX_TYPE( + GKO_DECLARE_BUILD_PARTITION_FROM_LOCAL_RANGE); + + +template +std::unique_ptr> +build_partition_from_local_size(std::shared_ptr exec, + mpi::communicator comm, size_type local_size) +{ + auto local_size_gi = static_cast(local_size); + array sizes(exec->get_master(), comm.size()); + comm.all_gather(exec, &local_size_gi, 1, sizes.get_data(), 1); + + array offsets(exec->get_master(), comm.size() + 1); + offsets.get_data()[0] = 0; + std::partial_sum(sizes.get_data(), sizes.get_data() + comm.size(), + offsets.get_data() + 1); + + return Partition::build_from_contiguous( + exec, offsets); +} + +#define GKO_DECLARE_BUILD_PARTITION_FROM_LOCAL_SIZE(_local_type, _global_type) \ + std::unique_ptr> \ + build_partition_from_local_size(std::shared_ptr exec, \ + mpi::communicator comm, \ + size_type local_range) +GKO_INSTANTIATE_FOR_EACH_LOCAL_GLOBAL_INDEX_TYPE( + GKO_DECLARE_BUILD_PARTITION_FROM_LOCAL_SIZE); + + +} // namespace distributed +} // namespace experimental +} // namespace gko diff --git a/core/distributed/partition_helpers_kernels.hpp b/core/distributed/partition_helpers_kernels.hpp new file mode 100644 index 00000000000..ed9fa60364f --- /dev/null +++ b/core/distributed/partition_helpers_kernels.hpp @@ -0,0 +1,86 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#ifndef GINKGO_PARTITION_HELPERS_KERNELS_HPP +#define GINKGO_PARTITION_HELPERS_KERNELS_HPP + + +#include + + +#include "core/base/kernel_declaration.hpp" + + +namespace gko { +namespace kernels { + + +#define GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START(_type) \ + void sort_by_range_start( \ + std::shared_ptr exec, \ + array<_type>& range_start_ends, \ + array& part_ids) + + +#define GKO_DECLARE_PARTITION_HELPERS_CHECK_CONSECUTIVE_RANGES(_type) \ + void check_consecutive_ranges(std::shared_ptr exec, \ + const array<_type>& range_start_ends, \ + bool& result) + + +#define GKO_DECLARE_PARTITION_HELPERS_COMPRESS_RANGES(_type) \ + void compress_ranges(std::shared_ptr exec, \ + const array<_type>& range_start_ends, \ + array<_type>& range_offsets) + + +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START(GlobalIndexType); \ + template \ + GKO_DECLARE_PARTITION_HELPERS_CHECK_CONSECUTIVE_RANGES(GlobalIndexType); \ + template \ + GKO_DECLARE_PARTITION_HELPERS_COMPRESS_RANGES(GlobalIndexType) + + +GKO_DECLARE_FOR_ALL_EXECUTOR_NAMESPACES(partition_helpers, + GKO_DECLARE_ALL_AS_TEMPLATES); + + +#undef GKO_DECLARE_ALL_AS_TEMPLATES + + +} // namespace kernels +} // namespace gko + + +#endif // GINKGO_PARTITION_HELPERS_KERNELS_HPP diff --git a/core/distributed/partition_kernels.hpp b/core/distributed/partition_kernels.hpp index 3d66ed113e8..070ff0839b4 100644 --- a/core/distributed/partition_kernels.hpp +++ b/core/distributed/partition_kernels.hpp @@ -49,10 +49,11 @@ namespace kernels { const array& mapping, \ size_type& num_ranges) -#define GKO_PARTITION_BUILD_FROM_CONTIGUOUS(GlobalIndexType) \ - void build_from_contiguous(std::shared_ptr exec, \ - const array& ranges, \ - GlobalIndexType* range_bounds, \ +#define GKO_PARTITION_BUILD_FROM_CONTIGUOUS(GlobalIndexType) \ + void build_from_contiguous(std::shared_ptr exec, \ + const array& ranges, \ + const array& part_id_mapping, \ + GlobalIndexType* range_bounds, \ comm_index_type* part_ids) #define GKO_PARTITION_BUILD_FROM_MAPPING(GlobalIndexType) \ diff --git a/core/test/base/iterator_factory.cpp b/core/test/base/iterator_factory.cpp index 68ed87e07cb..e4d8d39b340 100644 --- a/core/test/base/iterator_factory.cpp +++ b/core/test/base/iterator_factory.cpp @@ -67,13 +67,13 @@ namespace { template -class IteratorFactory : public ::testing::Test { +class ZipIterator : public ::testing::Test { protected: using value_type = typename std::tuple_element<0, decltype(ValueIndexType())>::type; using index_type = typename std::tuple_element<1, decltype(ValueIndexType())>::type; - IteratorFactory() + ZipIterator() : reversed_index{100, 50, 10, 9, 8, 7, 5, 5, 4, 3, 2, 1, 0, -1, -2}, ordered_index{-2, -1, 0, 1, 2, 3, 4, 5, 5, 7, 8, 9, 10, 50, 100}, reversed_value{15., 14., 13., 12., 11., 10., 9., 7., @@ -109,11 +109,11 @@ class IteratorFactory : public ::testing::Test { const std::vector ordered_value; }; -TYPED_TEST_SUITE(IteratorFactory, gko::test::ValueIndexTypes, +TYPED_TEST_SUITE(ZipIterator, gko::test::ValueIndexTypes, PairTypenameNameGenerator); -TYPED_TEST(IteratorFactory, EmptyIterator) +TYPED_TEST(ZipIterator, EmptyIterator) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -125,7 +125,7 @@ TYPED_TEST(IteratorFactory, EmptyIterator) } -TYPED_TEST(IteratorFactory, SortingReversedWithIterator) +TYPED_TEST(ZipIterator, SortingReversedWithIterator) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -140,7 +140,7 @@ TYPED_TEST(IteratorFactory, SortingReversedWithIterator) } -TYPED_TEST(IteratorFactory, SortingAlreadySortedWithIterator) +TYPED_TEST(ZipIterator, SortingAlreadySortedWithIterator) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -155,7 +155,7 @@ TYPED_TEST(IteratorFactory, SortingAlreadySortedWithIterator) } -TYPED_TEST(IteratorFactory, IteratorReferenceOperatorSmaller) +TYPED_TEST(ZipIterator, IteratorReferenceOperatorSmaller) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -170,7 +170,7 @@ TYPED_TEST(IteratorFactory, IteratorReferenceOperatorSmaller) } -TYPED_TEST(IteratorFactory, IteratorReferenceOperatorSmaller2) +TYPED_TEST(ZipIterator, IteratorReferenceOperatorSmaller2) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -185,7 +185,7 @@ TYPED_TEST(IteratorFactory, IteratorReferenceOperatorSmaller2) } -TYPED_TEST(IteratorFactory, IncreasingIterator) +TYPED_TEST(ZipIterator, IncreasingIterator) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -262,7 +262,7 @@ bool check_assertion_exit_code(int exit_code) } -TYPED_TEST(IteratorFactory, IncompatibleIteratorDeathTest) +TYPED_TEST(ZipIterator, IncompatibleIteratorDeathTest) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -286,7 +286,7 @@ TYPED_TEST(IteratorFactory, IncompatibleIteratorDeathTest) #endif -TYPED_TEST(IteratorFactory, DecreasingIterator) +TYPED_TEST(ZipIterator, DecreasingIterator) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -316,7 +316,7 @@ TYPED_TEST(IteratorFactory, DecreasingIterator) } -TYPED_TEST(IteratorFactory, CorrectDereferencing) +TYPED_TEST(ZipIterator, CorrectDereferencing) { using index_type_it = typename TestFixture::index_type; using value_type_it = typename TestFixture::value_type; @@ -337,7 +337,7 @@ TYPED_TEST(IteratorFactory, CorrectDereferencing) } -TYPED_TEST(IteratorFactory, CorrectSwapping) +TYPED_TEST(ZipIterator, CorrectSwapping) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -361,7 +361,7 @@ TYPED_TEST(IteratorFactory, CorrectSwapping) } -TYPED_TEST(IteratorFactory, CorrectHandWrittenSwapping) +TYPED_TEST(ZipIterator, CorrectHandWrittenSwapping) { using index_type = typename TestFixture::index_type; using value_type = typename TestFixture::value_type; @@ -388,4 +388,155 @@ TYPED_TEST(IteratorFactory, CorrectHandWrittenSwapping) } +template +class PermuteIterator : public ::testing::Test { +protected: + using value_type = ValueType; +}; + +TYPED_TEST_SUITE(PermuteIterator, gko::test::ValueAndIndexTypes, + TypenameNameGenerator); + + +TYPED_TEST(PermuteIterator, EmptyIterator) +{ + auto test_iter = gko::detail::make_permute_iterator( + nullptr, [](int i) { return i; }); + + ASSERT_NO_THROW(std::sort(test_iter, test_iter)); +} + + +TYPED_TEST(PermuteIterator, SortingWithIdentityPermutation) +{ + std::vector vec{6, 2, 5, 2, 4}; + std::vector sorted{2, 2, 4, 5, 6}; + auto test_iter = gko::detail::make_permute_iterator( + vec.begin(), [](int i) { return i; }); + + std::sort(test_iter, test_iter + vec.size()); + + ASSERT_EQ(vec, sorted); +} + + +TYPED_TEST(PermuteIterator, SortingWithReversePermutation) +{ + std::vector vec{6, 2, 5, 2, 4}; + std::vector sorted{6, 5, 4, 2, 2}; + auto test_iter = gko::detail::make_permute_iterator( + vec.begin(), + [size = vec.size()](int i) { return static_cast(size) - 1 - i; }); + + std::sort(test_iter, test_iter + vec.size()); + + ASSERT_EQ(vec, sorted); +} + + +TYPED_TEST(PermuteIterator, SortingWithStridedPermutation) +{ + std::vector vec{6, 8, 2, 9, 5, 1, 2, 7, 4, 0}; + std::vector sorted{2, 8, 2, 9, 4, 1, 5, 7, 6, 0}; + auto test_iter = gko::detail::make_permute_iterator( + vec.begin(), [](int i) { return 2 * i; }); + + std::sort(test_iter, test_iter + vec.size() / 2); + + ASSERT_EQ(vec, sorted); +} + + +TYPED_TEST(PermuteIterator, IncreasingIterator) +{ + std::vector vec{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + auto perm = [size = vec.size()](int i) { + return static_cast(size) - 1 - i; + }; + + auto test_iter = gko::detail::make_permute_iterator(vec.begin(), perm); + auto begin = test_iter; + auto plus_2 = begin + 2; + auto plus_2_rev = 2 + begin; + auto plus_minus_2 = plus_2 - 2; + auto increment_pre_2 = begin; + ++increment_pre_2; + ++increment_pre_2; + auto increment_post_2 = begin; + increment_post_2++; + increment_post_2++; + auto increment_pre_test = begin; + auto increment_post_test = begin; + + // check results for equality + ASSERT_TRUE(begin == plus_minus_2); + ASSERT_TRUE(plus_2 == increment_pre_2); + ASSERT_TRUE(plus_2_rev == increment_pre_2); + ASSERT_TRUE(increment_pre_2 == increment_post_2); + ASSERT_TRUE(begin == increment_post_test++); + ASSERT_TRUE(begin + 1 == ++increment_pre_test); + ASSERT_TRUE(*plus_2 == vec[perm(2)]); + // check other comparison operators and difference + std::vector::iterator, decltype(perm)>> + its{begin, + plus_2, + plus_2_rev, + plus_minus_2, + increment_pre_2, + increment_post_2, + increment_pre_test, + increment_post_test, + begin + 5, + begin + 9}; + std::sort(its.begin(), its.end()); + std::vector dists; + std::vector ref_dists{0, 1, 0, 1, 0, 0, 0, 3, 4}; + for (int i = 0; i < its.size() - 1; i++) { + SCOPED_TRACE(i); + dists.push_back(its[i + 1] - its[i]); + auto equal = dists.back() > 0; + ASSERT_EQ(its[i + 1] > its[i], equal); + ASSERT_EQ(its[i] < its[i + 1], equal); + ASSERT_EQ(its[i] != its[i + 1], equal); + ASSERT_EQ(its[i] == its[i + 1], !equal); + ASSERT_EQ(its[i] >= its[i + 1], !equal); + ASSERT_EQ(its[i + 1] <= its[i], !equal); + ASSERT_TRUE(its[i + 1] >= its[i]); + ASSERT_TRUE(its[i] <= its[i + 1]); + } + ASSERT_EQ(dists, ref_dists); +} + + +TYPED_TEST(PermuteIterator, DecreasingIterator) +{ + std::vector vec{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + auto perm = [size = vec.size()](int i) { + return static_cast(size) - 1 - i; + }; + + auto test_iter = gko::detail::make_permute_iterator(vec.begin(), perm); + + auto iter = test_iter + 5; + auto minus_2 = iter - 2; + auto minus_plus_2 = minus_2 + 2; + auto decrement_pre_2 = iter; + --decrement_pre_2; + --decrement_pre_2; + auto decrement_post_2 = iter; + decrement_post_2--; + decrement_post_2--; + auto decrement_pre_test = iter; + auto decrement_post_test = iter; + + ASSERT_TRUE(iter == minus_plus_2); + ASSERT_TRUE(minus_2 == decrement_pre_2); + ASSERT_TRUE(decrement_pre_2 == decrement_post_2); + ASSERT_TRUE(iter == decrement_post_test--); + ASSERT_TRUE(iter - 1 == --decrement_pre_test); + ASSERT_TRUE(*minus_2 == vec[perm(3)]); +} + + } // namespace diff --git a/cuda/CMakeLists.txt b/cuda/CMakeLists.txt index 764f47afb83..4c972d2a584 100644 --- a/cuda/CMakeLists.txt +++ b/cuda/CMakeLists.txt @@ -21,6 +21,7 @@ target_sources(ginkgo_cuda base/version.cpp components/prefix_sum_kernels.cu distributed/matrix_kernels.cu + distributed/partition_helpers_kernels.cu distributed/partition_kernels.cu distributed/vector_kernels.cu factorization/cholesky_kernels.cu diff --git a/cuda/distributed/partition_helpers_kernels.cu b/cuda/distributed/partition_helpers_kernels.cu new file mode 100644 index 00000000000..62dad1efaf1 --- /dev/null +++ b/cuda/distributed/partition_helpers_kernels.cu @@ -0,0 +1,57 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/distributed/partition_helpers_kernels.hpp" + + +#include +#include +#include +#include + + +#include "cuda/base/thrust.cuh" + + +namespace gko { +namespace kernels { +namespace cuda { +namespace partition_helpers { + + +#include "common/cuda_hip/distributed/partition_helpers_kernels.hpp.inc" + + +} // namespace partition_helpers +} // namespace cuda +} // namespace kernels +} // namespace gko diff --git a/dpcpp/CMakeLists.txt b/dpcpp/CMakeLists.txt index b70175c6b12..dd0d7c4cdfb 100644 --- a/dpcpp/CMakeLists.txt +++ b/dpcpp/CMakeLists.txt @@ -18,6 +18,7 @@ target_sources(ginkgo_dpcpp base/version.dp.cpp components/prefix_sum_kernels.dp.cpp distributed/matrix_kernels.dp.cpp + distributed/partition_helpers_kernels.dp.cpp distributed/partition_kernels.dp.cpp distributed/vector_kernels.dp.cpp factorization/cholesky_kernels.dp.cpp diff --git a/dpcpp/distributed/partition_helpers_kernels.dp.cpp b/dpcpp/distributed/partition_helpers_kernels.dp.cpp new file mode 100644 index 00000000000..8b0171cd349 --- /dev/null +++ b/dpcpp/distributed/partition_helpers_kernels.dp.cpp @@ -0,0 +1,119 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +// force-top: on +#include +#include +#include +// force-top: off + + +#include "core/distributed/partition_helpers_kernels.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +namespace partition_helpers { + +struct stride { +#if ONEDPL_VERSION_MAJOR >= 2022 && ONEDPL_VERSION_MINOR >= 1 + template + Index operator()(const Index& i) const + { + return i * 2; + } +#else + // Some older version require [] while some require (), so I added both + template + Index operator[](const Index& i) const + { + return i * 2; + } + + template + Index operator()(const Index& i) const + { + return i * 2; + } +#endif +}; + +template +void sort_by_range_start( + std::shared_ptr exec, + array& range_start_ends, + array& part_ids) +{ + auto policy = + oneapi::dpl::execution::make_device_policy(*exec->get_queue()); + auto num_ranges = range_start_ends.get_num_elems() / 2; + + auto start_it = oneapi::dpl::make_permutation_iterator( + range_start_ends.get_data(), stride{}); + auto end_it = oneapi::dpl::make_permutation_iterator( + range_start_ends.get_data() + 1, stride{}); + + // older versions of oneDPL have a bug when sorting permutation iterators +#if ONEDPL_VERSION_MAJOR >= 2022 && ONEDPL_VERSION_MINOR >= 1 + auto zip_it = + oneapi::dpl::make_zip_iterator(start_it, end_it, part_ids.get_data()); + std::stable_sort(policy, zip_it, zip_it + num_ranges, [](auto a, auto b) { + return std::get<0>(a) < std::get<0>(b); + }); +#else + array starts(exec, num_ranges); + array ends(exec, num_ranges); + + std::copy(policy, start_it, start_it + num_ranges, starts.get_data()); + std::copy(policy, end_it, end_it + num_ranges, ends.get_data()); + + auto zip_it = oneapi::dpl::make_zip_iterator( + starts.get_data(), ends.get_data(), part_ids.get_data()); + std::stable_sort(policy, zip_it, zip_it + num_ranges, [](auto a, auto b) { + return std::get<0>(a) < std::get<0>(b); + }); + + std::copy(policy, starts.get_data(), starts.get_data() + num_ranges, + start_it); + std::copy(policy, ends.get_data(), ends.get_data() + num_ranges, end_it); +#endif +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START); + + +} // namespace partition_helpers +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/hip/CMakeLists.txt b/hip/CMakeLists.txt index 5ec1718ca4d..779db13d36a 100644 --- a/hip/CMakeLists.txt +++ b/hip/CMakeLists.txt @@ -18,6 +18,7 @@ set(GINKGO_HIP_SOURCES base/version.hip.cpp components/prefix_sum_kernels.hip.cpp distributed/matrix_kernels.hip.cpp + distributed/partition_helpers_kernels.hip.cpp distributed/partition_kernels.hip.cpp distributed/vector_kernels.hip.cpp factorization/cholesky_kernels.hip.cpp diff --git a/hip/distributed/partition_helpers_kernels.hip.cpp b/hip/distributed/partition_helpers_kernels.hip.cpp new file mode 100644 index 00000000000..d4769141676 --- /dev/null +++ b/hip/distributed/partition_helpers_kernels.hip.cpp @@ -0,0 +1,57 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/distributed/partition_helpers_kernels.hpp" + + +#include +#include +#include +#include + + +#include "hip/base/thrust.hip.hpp" + + +namespace gko { +namespace kernels { +namespace hip { +namespace partition_helpers { + + +#include "common/cuda_hip/distributed/partition_helpers_kernels.hpp.inc" + + +} // namespace partition_helpers +} // namespace hip +} // namespace kernels +} // namespace gko diff --git a/include/ginkgo/core/distributed/partition.hpp b/include/ginkgo/core/distributed/partition.hpp index 0096edf999c..bb36528a4a8 100644 --- a/include/ginkgo/core/distributed/partition.hpp +++ b/include/ginkgo/core/distributed/partition.hpp @@ -83,11 +83,11 @@ namespace distributed { * ``` * starting_index[0] = 0, * starting_index[1] = 0, - * starting_index[2] = 3, // second range of part 1 + * starting_index[2] = 3, // second range of part 0 * starting_index[3] = 0, - * starting_index[4] = 5, // third range of part 1 + * starting_index[4] = 5, // third range of part 0 * ``` - * which you can use to iterate only over the the second range of part 1 (the + * which you can use to iterate only over the the second range of part 0 (the * third global range) with * ``` * for(int i = 0; i < r[3] - r[2]; ++i){ @@ -231,7 +231,7 @@ class Partition * * @return true if each part has no more than one contiguous range. */ - bool has_connected_parts(); + bool has_connected_parts() const; /** * Checks if the ranges are ordered by their part index. @@ -240,7 +240,7 @@ class Partition * * @return true if the ranges are ordered by their part index. */ - bool has_ordered_parts(); + bool has_ordered_parts() const; /** * Builds a partition from a given mapping global_index -> part_id. @@ -260,15 +260,18 @@ class Partition * * @param exec the Executor on which the partition should be built * @param ranges the boundaries of the ranges representing each part. - * Part i contains the indices [ranges[i], ranges[i + 1]). - * Has to contain at least one element. - * The first element has to be 0. + * Part part_id[i] contains the indices + * [ranges[i], ranges[i + 1]). Has to contain at least + * one element. The first element has to be 0. + * @param part_ids the part ids of the provided ranges. If empty, then + * it will assume range i belongs to part i. * * @return a Partition representing the given contiguous partitioning. */ static std::unique_ptr build_from_contiguous( std::shared_ptr exec, - const array& ranges); + const array& ranges, + const array& part_ids = {}); /** * Builds a partition by evenly distributing the global range. diff --git a/include/ginkgo/core/distributed/partition_helpers.hpp b/include/ginkgo/core/distributed/partition_helpers.hpp new file mode 100644 index 00000000000..6bc20350a7d --- /dev/null +++ b/include/ginkgo/core/distributed/partition_helpers.hpp @@ -0,0 +1,101 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#ifndef GKO_PUBLIC_CORE_DISTRIBUTED_PARTITION_HELPERS_HPP_ +#define GKO_PUBLIC_CORE_DISTRIBUTED_PARTITION_HELPERS_HPP_ + + +#include + + +#if GINKGO_BUILD_MPI + + +#include +#include + + +namespace gko { +namespace experimental { +namespace distributed { + +template +class Partition; + + +/** + * Builds a partition from a local range. + * + * @param exec the Executor on which the partition should be built. + * @param comm the communicator used to determine the global partition. + * @param local_range the start and end indices of the local range. + * + * @warning This throws, if the resulting partition would contain gaps. + * That means that for a partition of size `n` every local range `r_i + * = [s_i, e_i)` either `s_i != 0` and another local range `r_j = + * [s_j, e_j = s_i)` exists, or `e_i != n` and another local range + * `r_j = [s_j = e_i, e_j)` exists. + * + * @return a Partition where each range has the individual local_start + * and local_ends. + */ +template +std::unique_ptr> +build_partition_from_local_range(std::shared_ptr exec, + mpi::communicator comm, span local_range); + + +/** + * Builds a partition from a local size. + * + * @param exec the Executor on which the partition should be built. + * @param comm the communicator used to determine the global partition. + * @param local_range the number of the locally owned indices + * + * @return a Partition where each range has the specified local size. More + * specifically, if this is called on process i with local_size `s_i`, + * then the range `i` has size `s_i`, and range `r_i = [start, start + + * s_i)`, where `start = sum_j^(i-1) s_j`. + */ +template +std::unique_ptr> +build_partition_from_local_size(std::shared_ptr exec, + mpi::communicator comm, size_type local_size); + + +} // namespace distributed +} // namespace experimental +} // namespace gko + + +#endif // GINKGO_BUILD_MPI +#endif // GKO_PUBLIC_CORE_DISTRIBUTED_PARTITION_HELPERS_HPP_ diff --git a/include/ginkgo/ginkgo.hpp b/include/ginkgo/ginkgo.hpp index 179a8a01a46..594ad880b8c 100644 --- a/include/ginkgo/ginkgo.hpp +++ b/include/ginkgo/ginkgo.hpp @@ -82,6 +82,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include #include diff --git a/omp/CMakeLists.txt b/omp/CMakeLists.txt index 02248983385..c689ffc42f3 100644 --- a/omp/CMakeLists.txt +++ b/omp/CMakeLists.txt @@ -10,6 +10,7 @@ target_sources(ginkgo_omp base/version.cpp components/prefix_sum_kernels.cpp distributed/matrix_kernels.cpp + distributed/partition_helpers_kernels.cpp distributed/partition_kernels.cpp distributed/vector_kernels.cpp factorization/cholesky_kernels.cpp diff --git a/omp/distributed/partition_helpers_kernels.cpp b/omp/distributed/partition_helpers_kernels.cpp new file mode 100644 index 00000000000..2c006a22885 --- /dev/null +++ b/omp/distributed/partition_helpers_kernels.cpp @@ -0,0 +1,72 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/distributed/partition_helpers_kernels.hpp" + + +#include "core/base/iterator_factory.hpp" + + +namespace gko { +namespace kernels { +namespace omp { +namespace partition_helpers { + + +template +void sort_by_range_start( + std::shared_ptr exec, + array& range_start_ends, + array& part_ids) +{ + auto part_ids_d = part_ids.get_data(); + auto num_parts = part_ids.get_num_elems(); + auto start_it = detail::make_permute_iterator( + range_start_ends.get_data(), [](const auto i) { return 2 * i; }); + auto end_it = detail::make_permute_iterator( + range_start_ends.get_data() + 1, [](const auto i) { return 2 * i; }); + auto sort_it = detail::make_zip_iterator(start_it, end_it, part_ids_d); + // TODO: use TBB or parallel std with c++17 + std::stable_sort(sort_it, sort_it + num_parts, + [](const auto& a, const auto& b) { + return std::get<0>(a) < std::get<0>(b); + }); +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START); + + +} // namespace partition_helpers +} // namespace omp +} // namespace kernels +} // namespace gko diff --git a/reference/CMakeLists.txt b/reference/CMakeLists.txt index 074d5efe818..dd54e3fb52f 100644 --- a/reference/CMakeLists.txt +++ b/reference/CMakeLists.txt @@ -13,6 +13,7 @@ target_sources(ginkgo_reference components/precision_conversion_kernels.cpp components/prefix_sum_kernels.cpp distributed/matrix_kernels.cpp + distributed/partition_helpers_kernels.cpp distributed/partition_kernels.cpp distributed/vector_kernels.cpp factorization/cholesky_kernels.cpp diff --git a/reference/distributed/partition_helpers_kernels.cpp b/reference/distributed/partition_helpers_kernels.cpp new file mode 100644 index 00000000000..b68c10b1d01 --- /dev/null +++ b/reference/distributed/partition_helpers_kernels.cpp @@ -0,0 +1,114 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/distributed/partition_helpers_kernels.hpp" + + +#include "core/base/iterator_factory.hpp" + + +namespace gko { +namespace kernels { +namespace reference { +namespace partition_helpers { + + +template +void sort_by_range_start( + std::shared_ptr exec, + array& range_start_ends, + array& part_ids) +{ + auto part_ids_d = part_ids.get_data(); + auto num_parts = part_ids.get_num_elems(); + auto start_it = detail::make_permute_iterator( + range_start_ends.get_data(), [](const auto i) { return 2 * i; }); + auto end_it = detail::make_permute_iterator( + range_start_ends.get_data() + 1, [](const auto i) { return 2 * i; }); + auto sort_it = detail::make_zip_iterator(start_it, end_it, part_ids_d); + std::stable_sort(sort_it, sort_it + num_parts, + [](const auto& a, const auto& b) { + return std::get<0>(a) < std::get<0>(b); + }); +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_SORT_BY_RANGE_START); + + +template +void check_consecutive_ranges(std::shared_ptr exec, + const array& range_start_ends, + bool& result) +{ + auto num_parts = range_start_ends.get_num_elems() / 2; + auto start_it = + detail::make_permute_iterator(range_start_ends.get_const_data() + 2, + [](const auto i) { return 2 * i; }); + auto end_it = + detail::make_permute_iterator(range_start_ends.get_const_data() + 1, + [](const auto i) { return 2 * i; }); + auto range_it = detail::make_zip_iterator(start_it, end_it); + + if (num_parts) { + result = std::all_of( + range_it, range_it + num_parts - 1, + [](const auto& r) { return std::get<0>(r) == std::get<1>(r); }); + } else { + result = true; + } +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_CHECK_CONSECUTIVE_RANGES); + + +template +void compress_ranges(std::shared_ptr exec, + const array& range_start_ends, + array& range_offsets) +{ + range_offsets.get_data()[0] = range_start_ends.get_const_data()[0]; + for (int i = 0; i < range_offsets.get_num_elems() - 1; ++i) { + range_offsets.get_data()[i + 1] = + range_start_ends.get_const_data()[2 * i + 1]; + } +} + +GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE( + GKO_DECLARE_PARTITION_HELPERS_COMPRESS_RANGES); + + +} // namespace partition_helpers +} // namespace reference +} // namespace kernels +} // namespace gko diff --git a/reference/distributed/partition_kernels.cpp b/reference/distributed/partition_kernels.cpp index 6eae93d27d0..e9a2bfe7667 100644 --- a/reference/distributed/partition_kernels.cpp +++ b/reference/distributed/partition_kernels.cpp @@ -55,14 +55,16 @@ void count_ranges(std::shared_ptr exec, template void build_from_contiguous(std::shared_ptr exec, const array& ranges, + const array& part_id_mapping, GlobalIndexType* range_bounds, comm_index_type* part_ids) { + bool uses_mapping = part_id_mapping.get_num_elems() > 0; range_bounds[0] = 0; for (comm_index_type i = 0; i < ranges.get_num_elems() - 1; i++) { auto end = ranges.get_const_data()[i + 1]; range_bounds[i + 1] = end; - part_ids[i] = i; + part_ids[i] = uses_mapping ? part_id_mapping.get_const_data()[i] : i; } } diff --git a/reference/test/distributed/CMakeLists.txt b/reference/test/distributed/CMakeLists.txt index 2985c7b5e11..42ad2d7e1a2 100644 --- a/reference/test/distributed/CMakeLists.txt +++ b/reference/test/distributed/CMakeLists.txt @@ -1,3 +1,4 @@ ginkgo_create_test(matrix_kernels) +ginkgo_create_test(partition_helpers_kernels) ginkgo_create_test(partition_kernels) ginkgo_create_test(vector_kernels) diff --git a/reference/test/distributed/partition_helpers_kernels.cpp b/reference/test/distributed/partition_helpers_kernels.cpp new file mode 100644 index 00000000000..f0ce4918d01 --- /dev/null +++ b/reference/test/distributed/partition_helpers_kernels.cpp @@ -0,0 +1,145 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include +#include +#include + + +#include +#include + + +#include +#include + + +#include "core/distributed/partition_helpers_kernels.hpp" +#include "core/test/utils.hpp" + + +namespace { + + +using comm_index_type = gko::experimental::distributed::comm_index_type; + + +template +class PartitionHelpers : public ::testing::Test { +protected: + using global_index_type = GlobalIndexType; + + PartitionHelpers() : ref(gko::ReferenceExecutor::create()) {} + + std::shared_ptr ref; + gko::array default_range_start_ends{ + this->ref, {0, 4, 4, 7, 7, 9, 9, 11}}; + gko::array default_part_ids{this->ref, {0, 1, 2, 3}}; +}; + +TYPED_TEST_SUITE(PartitionHelpers, gko::test::IndexTypes, + TypenameNameGenerator); + + +TYPED_TEST(PartitionHelpers, CanSortByRangeStartIdentity) +{ + using itype = typename TestFixture::global_index_type; + auto range_start_ends = this->default_range_start_ends; + auto part_ids = this->default_part_ids; + + gko::kernels::reference::partition_helpers::sort_by_range_start( + this->ref, range_start_ends, part_ids); + + GKO_ASSERT_ARRAY_EQ(range_start_ends, this->default_range_start_ends); + GKO_ASSERT_ARRAY_EQ(part_ids, this->default_part_ids); +} + + +TYPED_TEST(PartitionHelpers, CanSortByRangeStart) +{ + using global_index_type = typename TestFixture::global_index_type; + gko::array range_start_ends{this->ref, + {7, 9, 4, 7, 0, 4, 9, 11}}; + gko::array result_part_ids{this->ref, {2, 1, 0, 3}}; + auto part_ids = this->default_part_ids; + + gko::kernels::reference::partition_helpers::sort_by_range_start( + this->ref, range_start_ends, part_ids); + + GKO_ASSERT_ARRAY_EQ(range_start_ends, this->default_range_start_ends); + GKO_ASSERT_ARRAY_EQ(part_ids, result_part_ids); +} + + +TYPED_TEST(PartitionHelpers, CanCheckConsecutiveRanges) +{ + using global_index_type = typename TestFixture::global_index_type; + auto range_start_ends = this->default_range_start_ends; + bool result = false; + + gko::kernels::reference::partition_helpers::check_consecutive_ranges( + this->ref, range_start_ends, result); + + ASSERT_TRUE(result); +} + + +TYPED_TEST(PartitionHelpers, CanCheckNonConsecutiveRanges) +{ + using global_index_type = typename TestFixture::global_index_type; + gko::array range_start_ends{this->ref, + {7, 9, 4, 7, 0, 4, 9, 11}}; + bool result = true; + + gko::kernels::reference::partition_helpers::check_consecutive_ranges( + this->ref, range_start_ends, result); + + ASSERT_FALSE(result); +} + + +TYPED_TEST(PartitionHelpers, CanCompressRanges) +{ + using itype = typename TestFixture::global_index_type; + auto range_start_ends = this->default_range_start_ends; + gko::array range_offsets{this->ref, + range_start_ends.get_num_elems() / 2 + 1}; + gko::array expected_range_offsets{this->ref, {0, 4, 7, 9, 11}}; + + gko::kernels::reference::partition_helpers::compress_ranges( + this->ref, range_start_ends, range_offsets); + + GKO_ASSERT_ARRAY_EQ(range_offsets, expected_range_offsets); +} + + +} // namespace diff --git a/reference/test/distributed/partition_kernels.cpp b/reference/test/distributed/partition_kernels.cpp index 4cc7750a193..f92349ee2eb 100644 --- a/reference/test/distributed/partition_kernels.cpp +++ b/reference/test/distributed/partition_kernels.cpp @@ -171,6 +171,28 @@ TYPED_TEST(Partition, BuildsFromRangeWithSingleElement) } +TYPED_TEST(Partition, BuildsFromRangesWithPartIds) +{ + using global_index_type = typename TestFixture::global_index_type; + using part_type = typename TestFixture::part_type; + gko::array ranges{this->ref, {0, 5, 5, 7, 9, 10}}; + gko::array part_id{this->ref, {0, 4, 3, 1, 2}}; + + auto partition = + part_type::build_from_contiguous(this->ref, ranges, part_id); + + EXPECT_EQ(partition->get_size(), + ranges.get_data()[ranges.get_num_elems() - 1]); + EXPECT_EQ(partition->get_num_ranges(), ranges.get_num_elems() - 1); + EXPECT_EQ(partition->get_num_parts(), ranges.get_num_elems() - 1); + EXPECT_EQ(partition->get_num_empty_parts(), 1); + assert_equal_data(partition->get_range_bounds(), {0, 5, 5, 7, 9, 10}); + assert_equal_data(partition->get_part_ids(), {0, 4, 3, 1, 2}); + assert_equal_data(partition->get_range_starting_indices(), {0, 0, 0, 0, 0}); + assert_equal_data(partition->get_part_sizes(), {5, 2, 1, 2, 0}); +} + + TYPED_TEST(Partition, BuildsFromGlobalSize) { using part_type = typename TestFixture::part_type; diff --git a/test/distributed/CMakeLists.txt b/test/distributed/CMakeLists.txt index 1c8e9b1e8fc..32b3810ea31 100644 --- a/test/distributed/CMakeLists.txt +++ b/test/distributed/CMakeLists.txt @@ -1,3 +1,4 @@ ginkgo_create_common_test(matrix_kernels DISABLE_EXECUTORS dpcpp) ginkgo_create_common_test(partition_kernels DISABLE_EXECUTORS dpcpp) ginkgo_create_common_test(vector_kernels DISABLE_EXECUTORS dpcpp) +ginkgo_create_common_and_reference_test(partition_helper_kernels) diff --git a/test/distributed/partition_helper_kernels.cpp b/test/distributed/partition_helper_kernels.cpp new file mode 100644 index 00000000000..a53505cf1f6 --- /dev/null +++ b/test/distributed/partition_helper_kernels.cpp @@ -0,0 +1,276 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include +#include + + +#include + + +#include "core/base/iterator_factory.hpp" +#include "core/distributed/partition_helpers_kernels.hpp" +#include "core/test/utils.hpp" +#include "test/utils/executor.hpp" + + +using gko::experimental::distributed::comm_index_type; + + +// TODO: remove with c++17 +template +T clamp(const T& v, const T& lo, const T& hi) +{ + return v < lo ? lo : (v > hi ? hi : v); +} + + +template +std::vector create_iota(IndexType min, IndexType max) +{ + std::vector iota(clamp(max - min, IndexType(0), max)); + std::iota(iota.begin(), iota.end(), min); + return iota; +} + + +template +std::vector create_range_offsets(gko::size_type num_ranges) +{ + std::default_random_engine engine; + std::uniform_int_distribution dist(5, 10); + std::vector range_sizes(num_ranges); + std::generate(range_sizes.begin(), range_sizes.end(), + [&]() { return dist(engine); }); + + std::vector range_offsets(num_ranges + 1, 0); + std::partial_sum(range_sizes.begin(), range_sizes.end(), + range_offsets.begin() + 1); + return range_offsets; +} + + +template +std::vector create_ranges( + const std::vector& range_offsets) +{ + assert(range_offsets.size() >= 2); + gko::size_type num_ranges = range_offsets.size() - 1; + std::vector ranges(num_ranges * 2, 0); + for (gko::size_type i = 1; i < num_ranges; ++i) { + ranges[2 * i - 1] = range_offsets[i]; + ranges[2 * i] = range_offsets[i]; + } + ranges.back() = range_offsets.back(); + return ranges; +} + + +template +std::vector create_ranges(gko::size_type num_ranges) +{ + auto range_offsets = create_range_offsets(num_ranges); + + return create_ranges(range_offsets); +} + + +std::vector sample_unique(std::size_t min, std::size_t max, + gko::size_type n) +{ + std::default_random_engine engine; + auto values = create_iota(min, max); + std::shuffle(values.begin(), values.end(), engine); + values.erase(values.begin() + clamp(n, gko::size_type(0), values.size()), + values.end()); + return values; +} + + +template +std::vector remove_indices(const std::vector& source, + std::vector idxs) +{ + std::sort(idxs.begin(), idxs.end(), std::greater<>{}); + auto result = source; + for (auto idx : idxs) { + result.erase(result.begin() + 2 * idx, result.begin() + 2 * idx + 1); + } + return result; +} + + +template +gko::array make_array(std::shared_ptr exec, + const std::vector& v) +{ + return gko::array(exec, v.begin(), v.end()); +} + + +template +std::pair, std::vector> +shuffle_range_and_pid(const std::vector& ranges, + const std::vector& pid) +{ + std::default_random_engine engine; + + auto result = std::make_pair(ranges, pid); + + auto num_ranges = result.second.size(); + auto range_start_it = gko::detail::make_permute_iterator( + result.first.begin(), [](const auto i) { return 2 * i; }); + auto range_end_it = gko::detail::make_permute_iterator( + result.first.begin() + 1, [](const auto i) { return 2 * i; }); + auto zip_it = gko::detail::make_zip_iterator(range_start_it, range_end_it, + result.second.begin()); + std::shuffle(zip_it, zip_it + num_ranges, engine); + + return result; +} + + +template +class PartitionHelpers : public CommonTestFixture { +protected: + using index_type = IndexType; +}; + +TYPED_TEST_SUITE(PartitionHelpers, gko::test::IndexTypes); + + +TYPED_TEST(PartitionHelpers, CanCheckConsecutiveRanges) +{ + using index_type = typename TestFixture::index_type; + auto offsets = make_array(this->exec, create_ranges(100)); + bool result = false; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::check_consecutive_ranges( + this->exec, offsets, result); + + ASSERT_TRUE(result); +} + + +TYPED_TEST(PartitionHelpers, CanCheckNonConsecutiveRanges) +{ + using index_type = typename TestFixture::index_type; + auto full_range_ends = create_ranges(100); + auto removal_idxs = sample_unique(0, full_range_ends.size() / 2, 4); + auto start_ends = + make_array(this->exec, remove_indices(full_range_ends, removal_idxs)); + bool result = true; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::check_consecutive_ranges( + this->exec, start_ends, result); + + ASSERT_FALSE(result); +} + + +TYPED_TEST(PartitionHelpers, CanCheckConsecutiveRangesWithSingleRange) +{ + using index_type = typename TestFixture::index_type; + auto start_ends = make_array(this->ref, create_ranges(1)); + bool result = false; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::check_consecutive_ranges( + this->exec, start_ends, result); + + ASSERT_TRUE(result); +} + + +TYPED_TEST(PartitionHelpers, CanCheckConsecutiveRangesWithSingleElement) +{ + using index_type = typename TestFixture::index_type; + auto start_ends = gko::array(this->exec, {1}); + bool result = false; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::check_consecutive_ranges( + this->exec, start_ends, result); + + ASSERT_TRUE(result); +} + + +TYPED_TEST(PartitionHelpers, CanSortConsecutiveRanges) +{ + using index_type = typename TestFixture::index_type; + auto start_ends = make_array(this->exec, create_ranges(100)); + auto part_ids = create_iota(0, 100); + auto part_ids_arr = gko::array( + this->exec, part_ids.begin(), part_ids.end()); + auto expected_start_ends = start_ends; + auto expected_part_ids = part_ids_arr; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::sort_by_range_start( + this->exec, start_ends, part_ids_arr); + + GKO_ASSERT_ARRAY_EQ(expected_start_ends, start_ends); + GKO_ASSERT_ARRAY_EQ(expected_part_ids, part_ids_arr); +} + + +TYPED_TEST(PartitionHelpers, CanSortNonConsecutiveRanges) +{ + using index_type = typename TestFixture::index_type; + auto ranges = create_ranges(100); + auto part_ids = create_iota(0, 100); + auto shuffled = shuffle_range_and_pid(ranges, part_ids); + auto expected_start_ends = make_array(this->exec, ranges); + auto expected_part_ids = gko::array( + this->exec, part_ids.begin(), part_ids.end()); + auto start_ends = make_array(this->exec, shuffled.first); + auto part_ids_arr = gko::array( + this->exec, shuffled.second.begin(), shuffled.second.end()); + + gko::kernels::EXEC_NAMESPACE::partition_helpers::sort_by_range_start( + this->exec, start_ends, part_ids_arr); + + GKO_ASSERT_ARRAY_EQ(expected_start_ends, start_ends); + GKO_ASSERT_ARRAY_EQ(expected_part_ids, part_ids_arr); +} + + +TYPED_TEST(PartitionHelpers, CanCompressRanges) +{ + using index_type = typename TestFixture::index_type; + auto expected_offsets = create_range_offsets(100); + auto ranges = make_array(this->exec, create_ranges(expected_offsets)); + gko::array offsets{this->exec, expected_offsets.size()}; + + gko::kernels::EXEC_NAMESPACE::partition_helpers::compress_ranges( + this->exec, ranges, offsets); + + GKO_ASSERT_ARRAY_EQ(offsets, make_array(this->exec, expected_offsets)); +} diff --git a/test/distributed/partition_kernels.cpp b/test/distributed/partition_kernels.cpp index 686d1432da5..7033abb37ef 100644 --- a/test/distributed/partition_kernels.cpp +++ b/test/distributed/partition_kernels.cpp @@ -276,6 +276,22 @@ TYPED_TEST(Partition, BuildsFromContiguousWithSingleEntry) } +TYPED_TEST(Partition, BuildsFromContiguousWithPartId) +{ + using global_index_type = typename TestFixture::global_index_type; + using part_type = typename TestFixture::part_type; + gko::array ranges{this->ref, + {0, 1234, 3134, 4578, 16435, 60000}}; + gko::array part_id{this->ref, {0, 4, 3, 1, 2}}; + gko::array dranges{this->exec, ranges}; + + auto part = part_type::build_from_contiguous(this->ref, ranges, part_id); + auto dpart = part_type::build_from_contiguous(this->exec, dranges, part_id); + + this->assert_equal(part, dpart); +} + + TYPED_TEST(Partition, BuildsFromGlobalSize) { using global_index_type = typename TestFixture::global_index_type; diff --git a/test/mpi/CMakeLists.txt b/test/mpi/CMakeLists.txt index 08050bde58f..fc0aec8138a 100644 --- a/test/mpi/CMakeLists.txt +++ b/test/mpi/CMakeLists.txt @@ -1,4 +1,5 @@ ginkgo_create_common_and_reference_test(matrix MPI_SIZE 3) +ginkgo_create_common_and_reference_test(partition_helpers MPI_SIZE 3) ginkgo_create_common_and_reference_test(vector MPI_SIZE 3) add_subdirectory(preconditioner) diff --git a/test/mpi/partition_helpers.cpp b/test/mpi/partition_helpers.cpp new file mode 100644 index 00000000000..de0b897fd13 --- /dev/null +++ b/test/mpi/partition_helpers.cpp @@ -0,0 +1,135 @@ +/************************************************************* +Copyright (c) 2017-2023, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include +#include + + +#include "core/test/utils.hpp" +#include "test/utils/mpi/executor.hpp" + + +using comm_index_type = gko::experimental::distributed::comm_index_type; + + +template +class PartitionHelpers : public CommonMpiTestFixture { +protected: + using index_type = IndexType; +}; + +TYPED_TEST_SUITE(PartitionHelpers, gko::test::IndexTypes, + TypenameNameGenerator); + + +TYPED_TEST(PartitionHelpers, CanBuildFromLocalRanges) +{ + using itype = typename TestFixture::index_type; + gko::span local_range[] = {{0u, 4u}, {4u, 9u}, {9u, 11u}}; + gko::array expects_ranges{this->exec, {0, 4, 9, 11}}; + gko::array expects_pid{this->exec, {0, 1, 2}}; + + auto part = + gko::experimental::distributed::build_partition_from_local_range< + gko::int32, itype>(this->exec, this->comm, + local_range[this->comm.rank()]); + + GKO_ASSERT_ARRAY_EQ( + expects_ranges, + gko::make_const_array_view(this->exec, expects_ranges.get_num_elems(), + part->get_range_bounds())); + GKO_ASSERT_ARRAY_EQ( + expects_pid, + gko::make_const_array_view(this->exec, expects_pid.get_num_elems(), + part->get_part_ids())); +} + + +TYPED_TEST(PartitionHelpers, CanBuildFromLocalRangesUnsorted) +{ + using itype = typename TestFixture::index_type; + gko::span local_range[] = {{4u, 9u}, {9u, 11u}, {0u, 4u}}; + gko::array expects_ranges{this->exec, {0, 4, 9, 11}}; + gko::array expects_pid{this->exec, {2, 0, 1}}; + + auto part = + gko::experimental::distributed::build_partition_from_local_range< + gko::int32, itype>(this->exec, this->comm, + local_range[this->comm.rank()]); + + GKO_ASSERT_ARRAY_EQ( + expects_ranges, + gko::make_const_array_view(this->exec, expects_ranges.get_num_elems(), + part->get_range_bounds())); + GKO_ASSERT_ARRAY_EQ( + expects_pid, + gko::make_const_array_view(this->exec, expects_pid.get_num_elems(), + part->get_part_ids())); +} + + +TYPED_TEST(PartitionHelpers, CanBuildFromLocalRangesThrowsOnGap) +{ + using itype = typename TestFixture::index_type; + gko::span local_range[] = {{4u, 6u}, {9u, 11u}, {0u, 4u}}; + // Hack because of multiple template arguments in macro + auto build_from_local_ranges = [](auto... args) { + return gko::experimental::distributed::build_partition_from_local_range< + gko::int32, itype>(args...); + }; + + ASSERT_THROW(build_from_local_ranges(this->exec, this->comm, + local_range[this->comm.rank()]), + gko::InvalidStateError); +} + + +TYPED_TEST(PartitionHelpers, CanBuildFromLocalSize) +{ + using itype = typename TestFixture::index_type; + gko::size_type local_range[] = {4, 5, 3}; + gko::array expects_ranges{this->exec, {0, 4, 9, 12}}; + gko::array expects_pid{this->exec, {0, 1, 2}}; + + auto part = gko::experimental::distributed::build_partition_from_local_size< + gko::int32, itype>(this->exec, this->comm, + local_range[this->comm.rank()]); + + GKO_ASSERT_ARRAY_EQ( + expects_ranges, + gko::make_const_array_view(this->exec, expects_ranges.get_num_elems(), + part->get_range_bounds())); + GKO_ASSERT_ARRAY_EQ( + expects_pid, + gko::make_const_array_view(this->exec, expects_pid.get_num_elems(), + part->get_part_ids())); +}