diff --git a/common/cuda_hip/preconditioner/batch_identity.hpp.inc b/common/cuda_hip/preconditioner/batch_identity.hpp.inc new file mode 100644 index 00000000000..4872e667112 --- /dev/null +++ b/common/cuda_hip/preconditioner/batch_identity.hpp.inc @@ -0,0 +1,89 @@ +/************************************************************* +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. +*************************************************************/ + +/** + * Identity preconditioner for batch solvers. ( To be able to have + * unpreconditioned solves ) + */ +template +class BatchIdentity final { +public: + using value_type = ValueType; + + /** + * The size of the work vector required in case of static allocation. + */ + static constexpr int work_size = 0; + + /** + * The size of the work vector required in case of dynamic allocation. + * + * For the Identity preconditioner, this is unnecessary, but this function + * is part of a 'batch preconditioner interface' because other + * preconditioners may need it. + */ + __host__ __device__ static constexpr int dynamic_work_size(int, int) + { + return 0; + } + + + /** + * Sets the input and generates the identity preconditioner.(Nothing needs + * to be actually generated.) + * + * @param mat Matrix for which to build an Ideniity preconditioner. + * @param work A 'work-vector', which is unneecessary here as no + * preconditioner values are to be stored. + */ + __device__ __forceinline__ void generate( + size_type, + const gko::batch::matrix::ell::batch_item& + mat, + ValueType*) + {} + + __device__ __forceinline__ void generate( + size_type, + const gko::batch::matrix::dense::batch_item& mat, + ValueType*) + {} + + __device__ __forceinline__ void apply(const int num_rows, + const ValueType* const r, + ValueType* const z) const + { + for (int li = threadIdx.x; li < num_rows; li += blockDim.x) { + z[li] = r[li]; + } + } +}; diff --git a/cuda/preconditioner/batch_preconditioners.cuh b/cuda/preconditioner/batch_preconditioners.cuh new file mode 100644 index 00000000000..aedfead3ef2 --- /dev/null +++ b/cuda/preconditioner/batch_preconditioners.cuh @@ -0,0 +1,58 @@ +/************************************************************* +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_CUDA_PRECONDITIONER_BATCH_PRECONDITIONERS_CUH_ +#define GKO_CUDA_PRECONDITIONER_BATCH_PRECONDITIONERS_CUH_ + + +#include + + +#include "core/matrix/batch_struct.hpp" +#include "cuda/components/cooperative_groups.cuh" +#include "cuda/components/load_store.cuh" +#include "cuda/components/reduction.cuh" + + +namespace gko { +namespace kernels { +namespace cuda { + + +#include "common/cuda_hip/preconditioner/batch_identity.hpp.inc" + + +} // namespace cuda +} // namespace kernels +} // namespace gko + +#endif // GKO_CUDA_PRECONDITIONER_BATCH_PRECONDITIONERS_CUH_ diff --git a/dpcpp/preconditioner/batch_identity.hpp.inc b/dpcpp/preconditioner/batch_identity.hpp.inc new file mode 100644 index 00000000000..fef9ac4fca3 --- /dev/null +++ b/dpcpp/preconditioner/batch_identity.hpp.inc @@ -0,0 +1,81 @@ +/************************************************************* +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. +*************************************************************/ + + +/** + * Identity preconditioner for batch solvers. ( To be able to have + * unpreconditioned solves ) + */ +template +class BatchIdentity final { +public: + using value_type = ValueType; + + /** + * The size of the work vector required in case of static allocation. + */ + static constexpr int work_size = 0; + + /** + * The size of the work vector required in case of dynamic allocation. + */ + static int dynamic_work_size(int, int) { return 0; } + + + /** + * Sets the input and generates the identity preconditioner.(Nothing needs + * to be actually generated.) + * + * @param mat Matrix for which to build an Ideniity preconditioner. + * @param work A 'work-vector', which is unneecessary here as no + * preconditioner values are to be stored. + */ + void generate(size_type batch_id, + const gko::batch::matrix::ell::batch_item&, + ValueType* const, sycl::nd_item<3> item_ct1) + {} + + void generate(size_type batch_id, + const gko::batch::matrix::dense::batch_item&, + ValueType* const, sycl::nd_item<3> item_ct1) + {} + + __dpct_inline__ void apply(const int num_rows, const ValueType* const r, + ValueType* const z, + sycl::nd_item<3> item_ct1) const + { + for (int li = item_ct1.get_local_linear_id(); li < num_rows; + li += item_ct1.get_local_range().size()) { + z[li] = r[li]; + } + } +}; diff --git a/dpcpp/preconditioner/batch_preconditioners.hpp b/dpcpp/preconditioner/batch_preconditioners.hpp new file mode 100644 index 00000000000..51865edb64a --- /dev/null +++ b/dpcpp/preconditioner/batch_preconditioners.hpp @@ -0,0 +1,55 @@ +/************************************************************* +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_DPCPP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_ +#define GKO_DPCPP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_ + + +#include + + +#include "core/matrix/batch_struct.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { + + +#include "dpcpp/preconditioner/batch_identity.hpp.inc" + + +} // namespace dpcpp +} // namespace kernels +} // namespace gko + +#endif // GKO_DPCPP_PRECONDITIONER_BATCH_PRECONDITIONERS_HPP_ diff --git a/hip/preconditioner/batch_preconditioners.hip.hpp b/hip/preconditioner/batch_preconditioners.hip.hpp new file mode 100644 index 00000000000..fdb36ad765c --- /dev/null +++ b/hip/preconditioner/batch_preconditioners.hip.hpp @@ -0,0 +1,59 @@ +/************************************************************* +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_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HIP_HPP_ +#define GKO_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HIP_HPP_ + + +#include + + +#include "core/matrix/batch_struct.hpp" +#include "hip/components/cooperative_groups.hip.hpp" +#include "hip/components/load_store.hip.hpp" +#include "hip/components/reduction.hip.hpp" + + +namespace gko { +namespace kernels { +namespace hip { + + +#include "common/cuda_hip/preconditioner/batch_identity.hpp.inc" + + +} // namespace hip +} // namespace kernels +} // namespace gko + + +#endif // GKO_HIP_PRECONDITIONER_BATCH_PRECONDITIONERS_HIP_HPP_ diff --git a/reference/preconditioner/batch_identity.hpp b/reference/preconditioner/batch_identity.hpp new file mode 100644 index 00000000000..c26beb155ac --- /dev/null +++ b/reference/preconditioner/batch_identity.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_REFERENCE_PRECONDITIONER_BATCH_IDENTITY_HPP_ +#define GKO_REFERENCE_PRECONDITIONER_BATCH_IDENTITY_HPP_ + + +#include "core/matrix/batch_struct.hpp" +#include "reference/base/config.hpp" + + +namespace gko { +namespace kernels { +namespace host { + + +/** + * Identity preconditioner for batch solvers. ( To be able to have + * unpreconditioned solves ) + */ +template +class BatchIdentity final { +public: + using value_type = ValueType; + + /** + * The size of the work vector required in case of static allocation. + */ + static constexpr int work_size = 0; + + /** + * The size of the work vector required in case of dynamic allocation. + */ + static int dynamic_work_size(int, int) { return 0; } + + + /** + * Sets the input and generates the identity preconditioner.(Nothing needs + * to be actually generated.) + * + * @param mat Matrix for which to build an Ideniity preconditioner. + * @param work A 'work-vector', which is unneecessary here as no + * preconditioner values are to be stored. + */ + void generate(size_type, + const gko::batch::matrix::ell::batch_item& mat, + ValueType* const work) + {} + + void generate( + size_type, + const gko::batch::matrix::dense::batch_item& mat, + ValueType* const work) + {} + + void apply(const gko::batch::multi_vector::batch_item& r, + const gko::batch::multi_vector::batch_item& z) const + { + for (int i = 0; i < r.num_rows; i++) { + for (int j = 0; j < r.num_rhs; j++) + z.values[i * z.stride + j] = r.values[i * r.stride + j]; + } + } +}; + + +} // namespace host +} // namespace kernels +} // namespace gko + +#endif // GKO_REFERENCE_PRECONDITIONER_BATCH_IDENTITY_HPP_