diff --git a/common/factorization/par_ilut_select_kernels.hpp.inc b/common/factorization/par_ilut_select_kernels.hpp.inc index 9b4897d1766..e443d7b6ba7 100644 --- a/common/factorization/par_ilut_select_kernels.hpp.inc +++ b/common/factorization/par_ilut_select_kernels.hpp.inc @@ -62,8 +62,7 @@ __global__ __launch_bounds__(searchtree_width) void build_searchtree( // assuming rounding towards zero auto stride = double(size) / sample_size; #pragma unroll - for (auto i = decltype(sampleselect_oversampling){0}; - i < sampleselect_oversampling; ++i) { + for (int i = 0; i < sampleselect_oversampling; ++i) { auto lidx = idx * sampleselect_oversampling + i; auto val = input[static_cast(lidx * stride)]; samples[i] = abs(val); @@ -120,8 +119,7 @@ __global__ __launch_bounds__(default_block_size) void count_buckets( auto el = abs(input[i]); IndexType tree_idx{}; #pragma unroll - for (auto level = decltype(sampleselect_searchtree_height){0}; - level < sampleselect_searchtree_height; ++level) { + for (int level = 0; level < sampleselect_searchtree_height; ++level) { auto cmp = !(el < sh_tree[tree_idx]); tree_idx = 2 * tree_idx + 1 + cmp; } @@ -170,7 +168,7 @@ __global__ __launch_bounds__(default_block_size) void block_prefix_sum( // compute prefix sum over warp-sized blocks IndexType total{}; auto base_idx = warp_idx * work_per_warp * warp.size(); - for (auto step = decltype(work_per_warp){0}; step < work_per_warp; ++step) { + for (IndexType step = 0; step < work_per_warp; ++step) { auto idx = warp_lane + step * warp.size() + base_idx; auto val = idx < num_blocks ? local_counters[idx] : zero(); IndexType warp_total{}; diff --git a/cuda/test/components/sorting_kernels.cu b/cuda/test/components/sorting_kernels.cu index f61bbd0694e..e2b7abc51d7 100644 --- a/cuda/test/components/sorting_kernels.cu +++ b/cuda/test/components/sorting_kernels.cu @@ -99,7 +99,7 @@ protected: { // we want some duplicate elements std::uniform_int_distribution dist(0, num_elements / 2); - for (auto i = decltype(num_elements){0}; i < num_elements; ++i) { + for (int i = 0; i < num_elements; ++i) { ref_shared.get_data()[i] = dist(rng); } ddata = gko::Array{cuda, ref_shared}; diff --git a/dpcpp/CMakeLists.txt b/dpcpp/CMakeLists.txt index 7729588d363..e2d476164e8 100644 --- a/dpcpp/CMakeLists.txt +++ b/dpcpp/CMakeLists.txt @@ -75,7 +75,7 @@ else () target_link_options(ginkgo_dpcpp PUBLIC -fsycl-device-code-split=per_kernel) endif() target_link_libraries(ginkgo_dpcpp PUBLIC ginkgo_device) -target_link_libraries(ginkgo_dpcpp PRIVATE $) +target_link_libraries(ginkgo_dpcpp PRIVATE MKL::MKL_DPCPP) if (GINKGO_DPCPP_SINGLE_MODE) target_compile_definitions(ginkgo_dpcpp PRIVATE GINKGO_DPCPP_SINGLE_MODE=1) endif() diff --git a/dpcpp/test/matrix/dense_kernels.cpp b/dpcpp/test/matrix/dense_kernels.cpp index 257ee6fbc6a..aada746c42a 100644 --- a/dpcpp/test/matrix/dense_kernels.cpp +++ b/dpcpp/test/matrix/dense_kernels.cpp @@ -51,6 +51,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/components/fill_array.hpp" #include "core/matrix/dense_kernels.hpp" #include "core/test/utils.hpp" +#include "dpcpp/test/utils.hpp" namespace { @@ -194,127 +195,6 @@ class Dense : public ::testing::Test { }; -TEST_F(Dense, DpcppFillIsEquivalentToRef) -{ - set_up_vector_data(3); - auto result = Mtx::create(ref); - - x->fill(42); - dx->fill(42); - result->copy_from(dx.get()); - - GKO_ASSERT_MTX_NEAR(result, x, r::value); -} - - -TEST_F(Dense, DpcppStridedFillIsEquivalentToRef) -{ - using T = vtype; - auto x = gko::initialize>( - 4, {I{1.0, 2.0}, I{3.0, 4.0}, I{5.0, 6.0}}, ref); - auto dx = gko::initialize>( - 4, {I{1.0, 2.0}, I{3.0, 4.0}, I{5.0, 6.0}}, dpcpp); - auto result = Mtx::create(ref); - - x->fill(42); - dx->fill(42); - result->copy_from(dx.get()); - - GKO_ASSERT_MTX_NEAR(result, x, r::value); -} - - -TEST_F(Dense, SingleVectorDpcppScaleIsEquivalentToRef) -{ - set_up_vector_data(1); - auto result = Mtx::create(ref); - - x->scale(alpha.get()); - dx->scale(dalpha.get()); - result->copy_from(dx.get()); - - GKO_ASSERT_MTX_NEAR(result, x, r::value); -} - - -TEST_F(Dense, MultipleVectorDpcppScaleIsEquivalentToRef) -{ - set_up_vector_data(20); - - x->scale(alpha.get()); - dx->scale(dalpha.get()); - - GKO_ASSERT_MTX_NEAR(dx, x, r::value); -} - - -TEST_F(Dense, MultipleVectorDpcppScaleWithDifferentAlphaIsEquivalentToRef) -{ - set_up_vector_data(20, true); - - x->scale(alpha.get()); - dx->scale(dalpha.get()); - - GKO_ASSERT_MTX_NEAR(dx, x, r::value); -} - - -TEST_F(Dense, SingleVectorDpcppAddScaledIsEquivalentToRef) -{ - set_up_vector_data(1); - - x->add_scaled(alpha.get(), y.get()); - dx->add_scaled(dalpha.get(), dy.get()); - - GKO_ASSERT_MTX_NEAR(dx, x, r::value); -} - - -TEST_F(Dense, MultipleVectorDpcppAddScaledIsEquivalentToRef) -{ - set_up_vector_data(20); - - x->add_scaled(alpha.get(), y.get()); - dx->add_scaled(dalpha.get(), dy.get()); - - GKO_ASSERT_MTX_NEAR(dx, x, r::value); -} - - -TEST_F(Dense, MultipleVectorDpcppAddScaledWithDifferentAlphaIsEquivalentToRef) -{ - set_up_vector_data(20); - - x->add_scaled(alpha.get(), y.get()); - dx->add_scaled(dalpha.get(), dy.get()); - - GKO_ASSERT_MTX_NEAR(dx, x, r::value); -} - - -TEST_F(Dense, AddsScaledDiagIsEquivalentToRef) -{ - auto mat = gen_mtx(532, 532); - gko::Array diag_values(ref, 532); - gko::kernels::reference::components::fill_array(ref, diag_values.get_data(), - 532, Mtx::value_type{2.0}); - auto diag = - gko::matrix::Diagonal::create(ref, 532, diag_values); - alpha = gko::initialize({2.0}, ref); - auto dmat = Mtx::create(dpcpp); - dmat->copy_from(mat.get()); - auto ddiag = gko::matrix::Diagonal::create(dpcpp); - ddiag->copy_from(diag.get()); - dalpha = Mtx::create(dpcpp); - dalpha->copy_from(alpha.get()); - - mat->add_scaled(alpha.get(), diag.get()); - dmat->add_scaled(dalpha.get(), ddiag.get()); - - GKO_ASSERT_MTX_NEAR(mat, dmat, r::value); -} - - TEST_F(Dense, SingleVectorDpcppComputeDotIsEquivalentToRef) { set_up_vector_data(1); @@ -384,11 +264,9 @@ TEST_F(Dense, SimpleApplyIsEquivalentToRef) } -#if !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, SimpleApplyMixedIsEquivalentToRef) { + skip_if_single_mode(); set_up_apply_data(); x->apply(convert(y).get(), convert(expected).get()); @@ -398,9 +276,6 @@ TEST_F(Dense, SimpleApplyMixedIsEquivalentToRef) } -#endif // !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, AdvancedApplyIsEquivalentToRef) { set_up_apply_data(); @@ -412,11 +287,9 @@ TEST_F(Dense, AdvancedApplyIsEquivalentToRef) } -#if !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, AdvancedApplyMixedIsEquivalentToRef) { + skip_if_single_mode(); set_up_apply_data(); x->apply(convert(alpha).get(), convert(y).get(), @@ -428,11 +301,9 @@ TEST_F(Dense, AdvancedApplyMixedIsEquivalentToRef) } -#endif // !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, ApplyToComplexIsEquivalentToRef) { + skip_if_single_mode(); set_up_apply_data(); auto complex_b = gen_mtx(25, 1); auto dcomplex_b = ComplexMtx::create(dpcpp); @@ -448,9 +319,6 @@ TEST_F(Dense, ApplyToComplexIsEquivalentToRef) } -#if !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, ApplyToMixedComplexIsEquivalentToRef) { set_up_apply_data(); @@ -467,8 +335,6 @@ TEST_F(Dense, ApplyToMixedComplexIsEquivalentToRef) GKO_ASSERT_MTX_NEAR(dcomplex_x, complex_x, 1e-7); } -#endif // !GINKGO_DPCPP_SINGLE_MODE - TEST_F(Dense, AdvancedApplyToComplexIsEquivalentToRef) { @@ -487,11 +353,9 @@ TEST_F(Dense, AdvancedApplyToComplexIsEquivalentToRef) } -#if !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, AdvancedApplyToMixedComplexIsEquivalentToRef) { + skip_if_single_mode(); set_up_apply_data(); auto complex_b = gen_mtx(25, 1); auto dcomplex_b = MixedComplexMtx::create(dpcpp); @@ -509,9 +373,6 @@ TEST_F(Dense, AdvancedApplyToMixedComplexIsEquivalentToRef) } -#endif // !GINKGO_DPCPP_SINGLE_MODE - - TEST_F(Dense, ComputeDotComplexIsEquivalentToRef) { set_up_apply_data(); @@ -733,219 +594,4 @@ TEST_F(Dense, CalculateTotalColsIsEquivalentToRef) } -TEST_F(Dense, CanGatherRows) -{ - set_up_apply_data(); - - auto r_gather = x->row_gather(rgather_idxs.get()); - auto dr_gather = dx->row_gather(rgather_idxs.get()); - - GKO_ASSERT_MTX_NEAR(r_gather.get(), dr_gather.get(), 0); -} - - -TEST_F(Dense, CanGatherRowsIntoDense) -{ - set_up_apply_data(); - auto gather_size = - gko::dim<2>{rgather_idxs->get_num_elems(), x->get_size()[1]}; - auto r_gather = Mtx::create(ref, gather_size); - // test make_temporary_clone and non-default stride - auto dr_gather = Mtx::create(ref, gather_size, x->get_size()[1] + 2); - - x->row_gather(rgather_idxs.get(), r_gather.get()); - dx->row_gather(rgather_idxs.get(), dr_gather.get()); - - GKO_ASSERT_MTX_NEAR(r_gather.get(), dr_gather.get(), 0); -} - - -TEST_F(Dense, IsPermutable) -{ - set_up_apply_data(); - - auto permuted = square->permute(rpermute_idxs.get()); - auto dpermuted = dsquare->permute(rpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(permuted.get()), - static_cast(dpermuted.get()), 0); -} - - -TEST_F(Dense, IsInversePermutable) -{ - set_up_apply_data(); - - auto permuted = square->inverse_permute(rpermute_idxs.get()); - auto dpermuted = dsquare->inverse_permute(rpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(permuted.get()), - static_cast(dpermuted.get()), 0); -} - - -TEST_F(Dense, IsRowPermutable) -{ - set_up_apply_data(); - - auto r_permute = x->row_permute(rpermute_idxs.get()); - auto dr_permute = dx->row_permute(rpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(r_permute.get()), - static_cast(dr_permute.get()), 0); -} - - -TEST_F(Dense, IsColPermutable) -{ - set_up_apply_data(); - - auto c_permute = x->column_permute(cpermute_idxs.get()); - auto dc_permute = dx->column_permute(cpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(c_permute.get()), - static_cast(dc_permute.get()), 0); -} - - -TEST_F(Dense, IsInverseRowPermutable) -{ - set_up_apply_data(); - - auto inverse_r_permute = x->inverse_row_permute(rpermute_idxs.get()); - auto d_inverse_r_permute = dx->inverse_row_permute(rpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(inverse_r_permute.get()), - static_cast(d_inverse_r_permute.get()), 0); -} - - -TEST_F(Dense, IsInverseColPermutable) -{ - set_up_apply_data(); - - auto inverse_c_permute = x->inverse_column_permute(cpermute_idxs.get()); - auto d_inverse_c_permute = dx->inverse_column_permute(cpermute_idxs.get()); - - GKO_ASSERT_MTX_NEAR(static_cast(inverse_c_permute.get()), - static_cast(d_inverse_c_permute.get()), 0); -} - - -TEST_F(Dense, ExtractDiagonalOnTallSkinnyIsEquivalentToRef) -{ - set_up_apply_data(); - - auto diag = x->extract_diagonal(); - auto ddiag = dx->extract_diagonal(); - - GKO_ASSERT_MTX_NEAR(diag.get(), ddiag.get(), 0); -} - - -TEST_F(Dense, ExtractDiagonalOnShortFatIsEquivalentToRef) -{ - set_up_apply_data(); - - auto diag = y->extract_diagonal(); - auto ddiag = dy->extract_diagonal(); - - GKO_ASSERT_MTX_NEAR(diag.get(), ddiag.get(), 0); -} - - -TEST_F(Dense, InplaceAbsoluteMatrixIsEquivalentToRef) -{ - set_up_apply_data(); - - x->compute_absolute_inplace(); - dx->compute_absolute_inplace(); - - GKO_ASSERT_MTX_NEAR(x, dx, r::value); -} - - -TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) -{ - set_up_apply_data(); - - auto abs_x = x->compute_absolute(); - auto dabs_x = dx->compute_absolute(); - - GKO_ASSERT_MTX_NEAR(abs_x, dabs_x, r::value); -} - - -TEST_F(Dense, MakeComplexIsEquivalentToRef) -{ - set_up_apply_data(); - - auto complex_x = x->make_complex(); - auto dcomplex_x = dx->make_complex(); - - GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); -} - - -TEST_F(Dense, MakeComplexWithGivenResultIsEquivalentToRef) -{ - set_up_apply_data(); - - auto complex_x = ComplexMtx::create(ref, x->get_size()); - x->make_complex(complex_x.get()); - auto dcomplex_x = ComplexMtx::create(dpcpp, x->get_size()); - dx->make_complex(dcomplex_x.get()); - - GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); -} - - -TEST_F(Dense, GetRealIsEquivalentToRef) -{ - set_up_apply_data(); - - auto real_x = x->get_real(); - auto dreal_x = dx->get_real(); - - GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); -} - - -TEST_F(Dense, GetRealWithGivenResultIsEquivalentToRef) -{ - set_up_apply_data(); - - auto real_x = Mtx::create(ref, x->get_size()); - x->get_real(real_x.get()); - auto dreal_x = Mtx::create(dpcpp, dx->get_size()); - dx->get_real(dreal_x.get()); - - GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); -} - - -TEST_F(Dense, GetImagIsEquivalentToRef) -{ - set_up_apply_data(); - - auto imag_x = x->get_imag(); - auto dimag_x = dx->get_imag(); - - GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); -} - - -TEST_F(Dense, GetImagWithGivenResultIsEquivalentToRef) -{ - set_up_apply_data(); - - auto imag_x = Mtx::create(ref, x->get_size()); - x->get_imag(imag_x.get()); - auto dimag_x = Mtx::create(dpcpp, dx->get_size()); - dx->get_imag(dimag_x.get()); - - GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); -} - - } // namespace diff --git a/dpcpp/test/utils.hpp b/dpcpp/test/utils.hpp new file mode 100644 index 00000000000..a11133d1bc6 --- /dev/null +++ b/dpcpp/test/utils.hpp @@ -0,0 +1,57 @@ +/************************************************************* +Copyright (c) 2017-2021, 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_TEST_UTILS_HPP_ +#define GKO_DPCPP_TEST_UTILS_HPP_ + + +#include + + +namespace { + + +/** + * skip the test if it is compiled in single mode + */ +void skip_if_single_mode() +{ +#if GINKGO_DPCPP_SINGLE_MODE + GTEST_SKIP(); +#endif // GINKGO_DPCPP_SINGLE_MODE +} + + +} // namespace + + +#endif // GKO_DPCPP_TEST_UTILS_HPP_