diff --git a/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc b/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc index efbbd323ef6..17a7e125332 100644 --- a/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc +++ b/common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc @@ -31,10 +31,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -/** - * Scales the vectors in global or shared memory with a factor of alpha (alpha - * is in global memory or shared memory) - */ template __device__ __forceinline__ void scale( const gko::batch_multi_vector::batch_entry& alpha, @@ -52,9 +48,9 @@ __device__ __forceinline__ void scale( template __global__ -__launch_bounds__(default_block_size, sm_multiplier) void scale_kernel( - const gko::batch_multi_vector::uniform_batch alpha, - const gko::batch_multi_vector::uniform_batch x, Mapping map) + __launch_bounds__(default_block_size, sm_multiplier) void scale_kernel( + const gko::batch_multi_vector::uniform_batch alpha, + const gko::batch_multi_vector::uniform_batch x, Mapping map) { for (size_type batch_id = blockIdx.x; batch_id < x.num_batch_entries; batch_id += gridDim.x) { @@ -83,10 +79,10 @@ __device__ __forceinline__ void add_scaled( template __global__ -__launch_bounds__(default_block_size, sm_multiplier) void add_scaled_kernel( - const gko::batch_multi_vector::uniform_batch alpha, - const gko::batch_multi_vector::uniform_batch x, - const gko::batch_multi_vector::uniform_batch y, Mapping map) + __launch_bounds__(default_block_size, sm_multiplier) void add_scaled_kernel( + const gko::batch_multi_vector::uniform_batch alpha, + const gko::batch_multi_vector::uniform_batch x, + const gko::batch_multi_vector::uniform_batch y, Mapping map) { for (size_type batch_id = blockIdx.x; batch_id < x.num_batch_entries; batch_id += gridDim.x) { @@ -222,11 +218,15 @@ __device__ __forceinline__ void compute_norm2( template -__global__ -__launch_bounds__(default_block_size, sm_multiplier) void compute_norm2_kernel( - const gko::batch_multi_vector::uniform_batch x, - const gko::batch_multi_vector::uniform_batch> - result) +__global__ __launch_bounds__( + default_block_size, + sm_multiplier) void compute_norm2_kernel(const gko::batch_multi_vector:: + uniform_batch + x, + const gko::batch_multi_vector:: + uniform_batch< + remove_complex> + result) { for (size_type batch_id = blockIdx.x; batch_id < x.num_batch_entries; batch_id += gridDim.x) { @@ -259,9 +259,9 @@ __device__ __forceinline__ void copy( template __global__ -__launch_bounds__(default_block_size, sm_multiplier) void copy_kernel( - const gko::batch_multi_vector::uniform_batch src, - const gko::batch_multi_vector::uniform_batch dst) + __launch_bounds__(default_block_size, sm_multiplier) void copy_kernel( + const gko::batch_multi_vector::uniform_batch src, + const gko::batch_multi_vector::uniform_batch dst) { for (size_type batch_id = blockIdx.x; batch_id < src.num_batch_entries; batch_id += gridDim.x) { diff --git a/core/base/batch_struct.hpp b/core/base/batch_struct.hpp index d85c413e691..ea1b3ef3f3f 100644 --- a/core/base/batch_struct.hpp +++ b/core/base/batch_struct.hpp @@ -44,21 +44,20 @@ namespace batch_multi_vector { /** - * Encapsulates one matrix from a batch of dense matrices (vectors). + * Encapsulates one matrix from a batch of multi-vectors. */ template struct batch_entry { using value_type = ValueType; ValueType* values; - size_type stride; + int stride; int num_rows; int num_rhs; }; + /** - * A 'simple' structure to store a global uniform batch of dense matrices. - * - * It is uniform in the sense that all matrices in the batch have common sizes. + * A 'simple' structure to store a global uniform batch of multi-vectors. */ template struct uniform_batch { @@ -67,7 +66,7 @@ struct uniform_batch { ValueType* values; size_type num_batch_entries; - size_type stride; + int stride; int num_rows; int num_rhs; @@ -122,8 +121,8 @@ batch_entry(const batch_multi_vector::uniform_batch& batch, template GKO_ATTRIBUTES GKO_INLINE batch_multi_vector::batch_entry -batch_entry(ValueType* const batch_values, const size_type stride, - const int num_rows, const int num_rhs, const size_type batch_idx) +batch_entry(ValueType* const batch_values, const int stride, const int num_rows, + const int num_rhs, const size_type batch_idx) { return {batch_values + batch_idx * stride * num_rows, stride, num_rows, num_rhs}; diff --git a/core/test/base/batch_multi_vector.cpp b/core/test/base/batch_multi_vector.cpp index 5fbc4d5aa32..e63ed883517 100644 --- a/core/test/base/batch_multi_vector.cpp +++ b/core/test/base/batch_multi_vector.cpp @@ -84,6 +84,7 @@ class BatchMultiVector : public ::testing::Test { { ASSERT_EQ(m->get_num_batch_entries(), 0); ASSERT_EQ(m->get_common_size(), gko::dim<2>{}); + ASSERT_EQ(m->get_const_values(), nullptr); } std::shared_ptr exec; @@ -100,13 +101,6 @@ TYPED_TEST(BatchMultiVector, CanBeEmpty) } -TYPED_TEST(BatchMultiVector, ReturnsNullValuesArrayWhenEmpty) -{ - auto empty = gko::BatchMultiVector::create(this->exec); - ASSERT_EQ(empty->get_const_values(), nullptr); -} - - TYPED_TEST(BatchMultiVector, KnowsItsSizeAndValues) { ASSERT_NE(this->mtx->get_const_values(), nullptr); @@ -165,10 +159,12 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromExistingData) using size_type = gko::size_type; // clang-format off value_type data[] = { - 1.0, 2.0, -1.0, - 3.0, 4.0, -1.0, - 3.0, 5.0, 1.0, - 5.0, 6.0, -3.0}; + 1.0, 2.0, + -1.0,3.0, + 4.0, -1.0, + 3.0, 5.0, + 1.0, 5.0, + 6.0, -3.0}; // clang-format on auto m = gko::BatchMultiVector::create( @@ -192,11 +188,13 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromExistingConstData) using value_type = typename TestFixture::value_type; using size_type = gko::size_type; // clang-format off - const value_type data[] = { - 1.0, 2.0, -1.0, - 3.0, 4.0, -1.0, - 3.0, 5.0, 1.0, - 5.0, 6.0, -3.0}; + value_type data[] = { + 1.0, 2.0, + -1.0,3.0, + 4.0, -1.0, + 3.0, 5.0, + 1.0, 5.0, + 6.0, -3.0}; // clang-format on auto m = gko::BatchMultiVector::create_const( @@ -215,7 +213,7 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromExistingConstData) } -TYPED_TEST(BatchMultiVector, CanBeConstructedFromBatchMultiVectorMatrices) +TYPED_TEST(BatchMultiVector, CanBeConstructedFromDenseMatrices) { using value_type = typename TestFixture::value_type; using DenseMtx = typename TestFixture::DenseMtx; @@ -227,12 +225,8 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromBatchMultiVectorMatrices) auto m = gko::BatchMultiVector::create( this->exec, std::vector{mat1.get(), mat2.get()}); - auto m_ref = gko::BatchMultiVector::create( - this->exec, std::vector{mat1.get(), mat2.get(), mat1.get(), - mat2.get(), mat1.get(), mat2.get()}); - auto m2 = gko::BatchMultiVector::create(this->exec, 3, m.get()); - GKO_ASSERT_BATCH_MTX_NEAR(m2.get(), m_ref.get(), 1e-14); + this->assert_equal_to_original_mtx(m.get()); } @@ -255,7 +249,7 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromDenseMatricesByDuplication) } -TYPED_TEST(BatchMultiVector, CanBeConstructedFromDenseMatrices) +TYPED_TEST(BatchMultiVector, CanBeConstructedFromBatchMultiVectorMatrices) { using value_type = typename TestFixture::value_type; using DenseMtx = typename TestFixture::DenseMtx; @@ -264,11 +258,15 @@ TYPED_TEST(BatchMultiVector, CanBeConstructedFromDenseMatrices) this->exec); auto mat2 = gko::initialize({{1.0, 2.5, 3.0}, {1.0, 2.0, 3.0}}, this->exec); - auto m = gko::BatchMultiVector::create( this->exec, std::vector{mat1.get(), mat2.get()}); + auto m_ref = gko::BatchMultiVector::create( + this->exec, std::vector{mat1.get(), mat2.get(), mat1.get(), + mat2.get(), mat1.get(), mat2.get()}); - this->assert_equal_to_original_mtx(m.get()); + auto m2 = gko::BatchMultiVector::create(this->exec, 3, m.get()); + + GKO_ASSERT_BATCH_MTX_NEAR(m2.get(), m_ref.get(), 1e-14); } @@ -356,6 +354,7 @@ TYPED_TEST(BatchMultiVector, CanBeUnbatchedIntoDenseMatrices) auto dense_mats = this->mtx->unbatch(); + ASSERT_EQ(dense_mats.size(), 2); GKO_ASSERT_MTX_NEAR(dense_mats[0].get(), mat1.get(), 0.); GKO_ASSERT_MTX_NEAR(dense_mats[1].get(), mat2.get(), 0.); } @@ -380,8 +379,8 @@ TYPED_TEST(BatchMultiVector, CanBeReadFromMatrixData) ASSERT_EQ(m->get_common_size(), gko::dim<2>(2, 2)); EXPECT_EQ(m->at(0, 0, 0), value_type{1.0}); - EXPECT_EQ(m->at(0, 1, 0), value_type{0.0}); EXPECT_EQ(m->at(0, 0, 1), value_type{3.0}); + EXPECT_EQ(m->at(0, 1, 0), value_type{0.0}); EXPECT_EQ(m->at(0, 1, 1), value_type{5.0}); EXPECT_EQ(m->at(1, 0, 0), value_type{-1.0}); EXPECT_EQ(m->at(1, 0, 1), value_type{0.5}); diff --git a/core/test/utils/assertions.hpp b/core/test/utils/assertions.hpp index 8e825a32d4f..44da77244f7 100644 --- a/core/test/utils/assertions.hpp +++ b/core/test/utils/assertions.hpp @@ -323,21 +323,18 @@ ::testing::AssertionResult batch_matrices_near_impl( const MatrixData2& second, double tolerance) { std::vector err; - std::vector err_flag; for (size_type b = 0; b < first.size(); ++b) { - auto num_rows = first[b].size[0]; - auto num_cols = first[b].size[1]; - if (num_rows != second[b].size[0] || num_cols != second[b].size[1]) { + if (first.size() != second.size()) { return ::testing::AssertionFailure() << "Expected matrices of equal size\n\t" << first_expression - << " is of size [" << num_rows << " x " << num_cols - << "]\n\t" << second_expression << " is of size [" - << second[b].size[0] << " x " << second[b].size[1] << "]" + << " is of size [" << first[b].size[0] << " x " + << first[b].size[1] << "]\n\t" << second_expression + << " is of size [" << second[b].size[0] << " x " + << second[b].size[1] << "]" << " for batch " << b; } err.push_back(detail::get_relative_error(first[b], second[b])); - err_flag.push_back(err.back() <= tolerance); } auto bat = std::find_if(err.begin(), err.end(), diff --git a/cuda/base/batch_multi_vector_kernels.cu b/cuda/base/batch_multi_vector_kernels.cu index 05e08be0adb..3fd80a2aa41 100644 --- a/cuda/base/batch_multi_vector_kernels.cu +++ b/cuda/base/batch_multi_vector_kernels.cu @@ -67,14 +67,16 @@ namespace batch_multi_vector { constexpr auto default_block_size = 256; constexpr int sm_multiplier = 4; +// clang-format off + // NOTE: DO NOT CHANGE THE ORDERING OF THE INCLUDES -// force-top: on + #include "common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc" -// force-top: off #include "common/cuda_hip/base/batch_multi_vector_kernel_launcher.hpp.inc" +// clang-format on } // namespace batch_multi_vector } // namespace cuda diff --git a/cuda/base/batch_struct.hpp b/cuda/base/batch_struct.hpp index f9a50376362..d9907b41531 100644 --- a/cuda/base/batch_struct.hpp +++ b/cuda/base/batch_struct.hpp @@ -66,7 +66,7 @@ inline gko::batch_multi_vector::uniform_batch> get_batch_struct(const BatchMultiVector* const op) { return {as_cuda_type(op->get_const_values()), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } @@ -79,7 +79,7 @@ inline gko::batch_multi_vector::uniform_batch> get_batch_struct(BatchMultiVector* const op) { return {as_cuda_type(op->get_values()), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } diff --git a/dpcpp/base/batch_multi_vector_kernels.dp.cpp b/dpcpp/base/batch_multi_vector_kernels.dp.cpp index 2c48970d13d..e27b3fc810f 100644 --- a/dpcpp/base/batch_multi_vector_kernels.dp.cpp +++ b/dpcpp/base/batch_multi_vector_kernels.dp.cpp @@ -86,7 +86,7 @@ void scale(std::shared_ptr exec, // Launch a kernel that has nbatches blocks, each block has max group size if (alpha->get_common_size()[1] == 1) { - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) { auto group = item_ct1.get_group(); @@ -98,7 +98,7 @@ void scale(std::shared_ptr exec, }); }); } else { - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) { auto group = item_ct1.get_group(); @@ -136,7 +136,7 @@ void add_scaled(std::shared_ptr exec, const auto x_ub = get_batch_struct(x); const auto y_ub = get_batch_struct(y); if (alpha->get_common_size()[1] == 1) { - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) { auto group = item_ct1.get_group(); @@ -149,7 +149,7 @@ void add_scaled(std::shared_ptr exec, }); }); } else { - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) { auto group = item_ct1.get_group(); @@ -187,7 +187,7 @@ void compute_dot(std::shared_ptr exec, const dim3 grid(num_batches); // TODO: Remove reqd_sub_group size and use sycl::reduce_over_group - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [= ](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size( @@ -225,7 +225,7 @@ void compute_conj_dot(std::shared_ptr exec, const dim3 block(group_size); const dim3 grid(num_batches); - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [= ](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size( @@ -262,7 +262,7 @@ void compute_norm2(std::shared_ptr exec, const dim3 block(group_size); const dim3 grid(num_batches); - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [= ](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size( @@ -296,7 +296,7 @@ void copy(std::shared_ptr exec, const dim3 block(group_size); const dim3 grid(num_batches); - (exec->get_queue())->submit([&](sycl::handler& cgh) { + exec->get_queue()->submit([&](sycl::handler& cgh) { cgh.parallel_for( sycl_nd_range(grid, block), [=](sycl::nd_item<3> item_ct1) { auto group = item_ct1.get_group(); diff --git a/dpcpp/base/batch_struct.hpp b/dpcpp/base/batch_struct.hpp index 1a83fad020c..c9ee5800b3e 100644 --- a/dpcpp/base/batch_struct.hpp +++ b/dpcpp/base/batch_struct.hpp @@ -65,7 +65,7 @@ inline gko::batch_multi_vector::uniform_batch get_batch_struct( const BatchMultiVector* const op) { return {op->get_const_values(), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } @@ -79,7 +79,7 @@ inline gko::batch_multi_vector::uniform_batch get_batch_struct( BatchMultiVector* const op) { return {op->get_values(), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } diff --git a/hip/base/batch_multi_vector_kernels.hip.cpp b/hip/base/batch_multi_vector_kernels.hip.cpp index 096c5e8a5d3..40e828b5d45 100644 --- a/hip/base/batch_multi_vector_kernels.hip.cpp +++ b/hip/base/batch_multi_vector_kernels.hip.cpp @@ -69,14 +69,17 @@ constexpr auto default_block_size = 256; constexpr int sm_multiplier = 4; +// clang-format off + // NOTE: DO NOT CHANGE THE ORDERING OF THE INCLUDES -// force-top: on + #include "common/cuda_hip/base/batch_multi_vector_kernels.hpp.inc" -// force-top: off #include "common/cuda_hip/base/batch_multi_vector_kernel_launcher.hpp.inc" +// clang-format on + } // namespace batch_multi_vector } // namespace hip diff --git a/hip/base/batch_struct.hip.hpp b/hip/base/batch_struct.hip.hpp index bff659838bd..3171e7e1df8 100644 --- a/hip/base/batch_struct.hip.hpp +++ b/hip/base/batch_struct.hip.hpp @@ -66,7 +66,7 @@ inline gko::batch_multi_vector::uniform_batch> get_batch_struct(const BatchMultiVector* const op) { return {as_hip_type(op->get_const_values()), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } @@ -79,7 +79,7 @@ inline gko::batch_multi_vector::uniform_batch> get_batch_struct(BatchMultiVector* const op) { return {as_hip_type(op->get_values()), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } diff --git a/include/ginkgo/core/base/batch_dim.hpp b/include/ginkgo/core/base/batch_dim.hpp index 37ce5993220..c52b732f610 100644 --- a/include/ginkgo/core/base/batch_dim.hpp +++ b/include/ginkgo/core/base/batch_dim.hpp @@ -102,7 +102,7 @@ struct batch_dim { /** - * Checks if two batch dim objects are different. + * Checks if two batch_dim objects are different. * * @tparam Dimensionality number of dimensions of the dim objects * @tparam DimensionType datatype used to represent each dimension @@ -123,7 +123,7 @@ struct batch_dim { * Creates a batch_dim object which stores a uniform size for all batch * entries. * - * @param num_batch_entries number of batch entries to be stored + * @param num_batch_entries the number of batch entries to be stored * @param common_size the common size of all the batch entries stored * * @note Use this constructor when uniform batches need to be stored. diff --git a/include/ginkgo/core/base/batch_multi_vector.hpp b/include/ginkgo/core/base/batch_multi_vector.hpp index ac4a2feb419..b91c50966a1 100644 --- a/include/ginkgo/core/base/batch_multi_vector.hpp +++ b/include/ginkgo/core/base/batch_multi_vector.hpp @@ -170,26 +170,6 @@ class BatchMultiVector */ dim<2> get_common_size() const { return batch_size_.get_common_size(); } - /** - * Returns a pointer to the array of values of the beginning of the batched - * multi-vector. - * - * @return the pointer to the array of values - */ - value_type* get_values() noexcept { return values_.get_data(); } - - /** - * @copydoc get_values() - * - * @note This is the constant version of the function, which can be - * significantly more memory efficient than the non-constant version, - * so always prefer this version. - */ - const value_type* get_const_values() const noexcept - { - return values_.get_const_data(); - } - /** * Returns a pointer to the array of values of the multi-vector for a * specific batch entry. @@ -198,7 +178,7 @@ class BatchMultiVector * * @return the pointer to the array of values */ - value_type* get_values(size_type batch_id) noexcept + value_type* get_values(size_type batch_id = 0) noexcept { GKO_ASSERT(batch_id < this->get_num_batch_entries()); return values_.get_data() + @@ -212,7 +192,7 @@ class BatchMultiVector * significantly more memory efficient than the non-constant version, * so always prefer this version. */ - const value_type* get_const_values(size_type batch_id) const noexcept + const value_type* get_const_values(size_type batch_id = 0) const noexcept { GKO_ASSERT(batch_id < this->get_num_batch_entries()); return values_.get_const_data() + diff --git a/reference/base/batch_struct.hpp b/reference/base/batch_struct.hpp index ed1350dc366..41262be1d48 100644 --- a/reference/base/batch_struct.hpp +++ b/reference/base/batch_struct.hpp @@ -67,7 +67,7 @@ inline gko::batch_multi_vector::uniform_batch get_batch_struct( const BatchMultiVector* const op) { return {op->get_const_values(), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } @@ -81,7 +81,7 @@ inline gko::batch_multi_vector::uniform_batch get_batch_struct( BatchMultiVector* const op) { return {op->get_values(), op->get_num_batch_entries(), - op->get_common_size()[1], + static_cast(op->get_common_size()[1]), static_cast(op->get_common_size()[0]), static_cast(op->get_common_size()[1])}; } diff --git a/reference/test/base/batch_multi_vector_kernels.cpp b/reference/test/base/batch_multi_vector_kernels.cpp index 445cdedb73f..f6ae66d8249 100644 --- a/reference/test/base/batch_multi_vector_kernels.cpp +++ b/reference/test/base/batch_multi_vector_kernels.cpp @@ -127,7 +127,6 @@ class BatchMultiVector : public ::testing::Test { std::ranlux48 rand_engine; }; - TYPED_TEST_SUITE(BatchMultiVector, gko::test::ValueTypes); @@ -137,7 +136,6 @@ TYPED_TEST(BatchMultiVector, ScalesData) using T = typename TestFixture::value_type; auto alpha = gko::batch_initialize( {{{2.0, -2.0, 1.5}}, {{3.0, -1.0, 0.25}}}, this->exec); - auto ualpha = alpha->unbatch(); this->mtx_0->scale(alpha.get()); @@ -155,7 +153,6 @@ TYPED_TEST(BatchMultiVector, ScalesDataWithScalar) using Mtx = typename TestFixture::Mtx; using T = typename TestFixture::value_type; auto alpha = gko::batch_initialize({{2.0}, {-2.0}}, this->exec); - auto ualpha = alpha->unbatch(); this->mtx_1->scale(alpha.get()); @@ -174,7 +171,6 @@ TYPED_TEST(BatchMultiVector, ScalesDataWithStride) using T = typename TestFixture::value_type; auto alpha = gko::batch_initialize( {{{2.0, -2.0, -1.5}}, {{2.0, -2.0, 3.0}}}, this->exec); - auto ualpha = alpha->unbatch(); this->mtx_1->scale(alpha.get()); @@ -193,7 +189,6 @@ TYPED_TEST(BatchMultiVector, AddsScaled) using T = typename TestFixture::value_type; auto alpha = gko::batch_initialize( {{{2.0, -2.0, 1.5}}, {{2.0, -2.0, 3.0}}}, this->exec); - auto ualpha = alpha->unbatch(); this->mtx_1->add_scaled(alpha.get(), this->mtx_0.get()); @@ -211,7 +206,6 @@ TYPED_TEST(BatchMultiVector, AddsScaledWithScalar) using Mtx = typename TestFixture::Mtx; using T = typename TestFixture::value_type; auto alpha = gko::batch_initialize({{2.0}, {-2.0}}, this->exec); - auto ualpha = alpha->unbatch(); this->mtx_1->add_scaled(alpha.get(), this->mtx_0.get()); @@ -241,7 +235,6 @@ TYPED_TEST(BatchMultiVector, ComputesDot) using T = typename TestFixture::value_type; auto result = Mtx::create(this->exec, gko::batch_dim<2>(2, gko::dim<2>{1, 3})); - auto ures = result->unbatch(); this->mtx_0->compute_dot(this->mtx_1.get(), result.get()); @@ -286,7 +279,6 @@ TYPED_TEST(BatchMultiVector, ComputesConjDot) using T = typename TestFixture::value_type; auto result = Mtx::create(this->exec, gko::batch_dim<2>(2, gko::dim<2>{1, 3})); - auto ures = result->unbatch(); this->mtx_0->compute_conj_dot(this->mtx_1.get(), result.get()); diff --git a/test/base/batch_multi_vector_kernels.cpp b/test/base/batch_multi_vector_kernels.cpp index 631b9a10c24..015adbce798 100644 --- a/test/base/batch_multi_vector_kernels.cpp +++ b/test/base/batch_multi_vector_kernels.cpp @@ -82,46 +82,16 @@ class BatchMultiVector : public CommonTestFixture { alpha = gko::batch_initialize(batch_size, {2.0}, ref); beta = gko::batch_initialize(batch_size, {-0.5}, ref); } - dx = Mtx::create(exec); - dx->copy_from(x.get()); - dy = Mtx::create(exec); - dy->copy_from(y.get()); - dalpha = Mtx::create(exec); - dalpha->copy_from(alpha.get()); - dbeta = gko::clone(exec, beta.get()); + dx = gko::clone(exec, x); + dy = gko::clone(exec, y); + dalpha = gko::clone(exec, alpha); + dbeta = gko::clone(exec, beta); expected = Mtx::create( ref, gko::batch_dim<2>(batch_size, gko::dim<2>{1, num_vecs})); dresult = Mtx::create( exec, gko::batch_dim<2>(batch_size, gko::dim<2>{1, num_vecs})); } - void set_up_apply_data(const int p = 1) - { - const int m = 35, n = 15; - x = gen_mtx(batch_size, m, n); - c_x = gen_mtx(batch_size, m, n); - y = gen_mtx(batch_size, n, p); - expected = gen_mtx(batch_size, m, p); - alpha = gko::batch_initialize(batch_size, {2.0}, ref); - beta = gko::batch_initialize(batch_size, {-1.0}, ref); - square = gen_mtx(batch_size, x->get_common_size()[0], - x->get_common_size()[0]); - dx = Mtx::create(exec); - dx->copy_from(x.get()); - dc_x = ComplexMtx::create(exec); - dc_x->copy_from(c_x.get()); - dy = Mtx::create(exec); - dy->copy_from(y.get()); - dresult = Mtx::create(exec); - dresult->copy_from(expected.get()); - dalpha = Mtx::create(exec); - dalpha->copy_from(alpha.get()); - dbeta = Mtx::create(exec); - dbeta->copy_from(beta.get()); - dsquare = Mtx::create(exec); - dsquare->copy_from(square.get()); - } - std::ranlux48 rand_engine; const size_t batch_size = 11; diff --git a/test/test_install/test_install.cpp b/test/test_install/test_install.cpp index 2016f00dade..2467e99f62b 100644 --- a/test/test_install/test_install.cpp +++ b/test/test_install/test_install.cpp @@ -210,6 +210,20 @@ int main() array_type test; } + // core/base/batch_dim.hpp + { + using type1 = int; + auto common_size = gko::dim<2>{4, 2}; + auto test = gko::batch_dim<2, type1>{2, common_size}; + } + + // core/base/batch_multi_vector.hpp + { + using type1 = float; + using batch_multi_vector_type = gko::BatchMultiVector; + auto test = batch_multi_vector_type::create(exec); + } + // core/base/combination.hpp { using type1 = int;