Dia/Cuda/Morpheus_Multiply_Impl.hpp

 
#ifndef MORPHEUS_DIA_CUDA_MULTIPLY_IMPL_HPP
#define MORPHEUS_DIA_CUDA_MULTIPLY_IMPL_HPP
 
#include <Morpheus_Macros.hpp>
#if defined(MORPHEUS_ENABLE_CUDA)
 
#include <Morpheus_SpaceTraits.hpp>
#include <Morpheus_FormatTraits.hpp>
#include <Morpheus_FormatTags.hpp>
#include <Morpheus_Spaces.hpp>
 
#include <impl/Morpheus_CudaUtils.hpp>
#include <impl/Dia/Kernels/Morpheus_Multiply_Impl.hpp>
 
namespace Morpheus {
namespace Impl {
 
template <typename ExecSpace, typename Matrix, typename Vector>
inline void multiply(
    const Matrix& A, const Vector& x, Vector& y, const bool init,
    typename std::enable_if_t<
        Morpheus::is_dia_matrix_format_container_v<Matrix> &&
        Morpheus::is_dense_vector_format_container_v<Vector> &&
        Morpheus::has_custom_backend_v<ExecSpace> &&
        Morpheus::has_cuda_execution_space_v<ExecSpace> &&
        Morpheus::has_access_v<ExecSpace, Matrix, Vector>>* = nullptr) {
  using index_type = typename Matrix::index_type;
  using size_type  = typename Matrix::size_type;
  using value_type = typename Matrix::value_type;
 
  const size_type BLOCK_SIZE = 256;
  const size_type MAX_BLOCKS = max_active_blocks(
      Kernels::spmv_dia_kernel<size_type, index_type, value_type, BLOCK_SIZE>,
      BLOCK_SIZE, (size_type)sizeof(index_type) * BLOCK_SIZE);
  const size_type NUM_BLOCKS =
      std::min<size_type>(MAX_BLOCKS, DIVIDE_INTO(A.nrows(), BLOCK_SIZE));
 
  const index_type* D     = A.cdiagonal_offsets().data();
  const value_type* V     = A.cvalues().data();
  const value_type* x_ptr = x.data();
  value_type* y_ptr       = y.data();
 
  const index_type num_diagonals = A.cvalues().ncols();
  const index_type pitch         = A.cvalues().nrows();
 
  if (num_diagonals == 0) {
    // empty matrix
    return;
  }
 
  if (init) {
    y.assign(y.size(), 0);
  }
 
  Kernels::spmv_dia_kernel<size_type, index_type, value_type, BLOCK_SIZE>
      <<<NUM_BLOCKS, BLOCK_SIZE, 0>>>(A.nrows(), A.ncols(), num_diagonals,
                                      pitch, D, V, x_ptr, y_ptr);
 
#if defined(DEBUG) || defined(MORPHEUS_DEBUG)
  getLastCudaError("spmv_dia_kernel: Kernel execution failed");
#endif
}
 
}  // namespace Impl
}  // namespace Morpheus
 
#endif  // MORPHEUS_ENABLE_CUDA
#endif  // MORPHEUS_DIA_CUDA_MULTIPLY_IMPL_HPP