From 281447e52167d496dee14bab3037be3e0b7e2a17 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Thomas=20Gr=C3=BCtzmacher?= <thomas.gruetzmacher@kit.edu>
Date: Wed, 17 Mar 2021 11:29:28 +0100
Subject: [PATCH] Fix HIP and CUDA ELL SpMV Debug error

---
 cuda/matrix/ell_kernels.cu     | 4 ++--
 hip/matrix/ell_kernels.hip.cpp | 4 ++--
 2 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/cuda/matrix/ell_kernels.cu b/cuda/matrix/ell_kernels.cu
index 2dc669ba5a7..57617269ed0 100644
--- a/cuda/matrix/ell_kernels.cu
+++ b/cuda/matrix/ell_kernels.cu
@@ -151,8 +151,8 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
         std::array<size_type, 1>{{num_stored_elements_per_row * stride}},
         a->get_const_values());
     const auto b_vals = gko::acc::range<b_accessor>(
-        std::array<size_type, 2>{{nrows, b->get_stride()}},
-        b->get_const_values());
+        std::array<size_type, 2>{{b->get_size()[0], b->get_size()[1]}},
+        b->get_const_values(), std::array<size_type, 1>{{b->get_stride()}});
 
     if (alpha == nullptr && beta == nullptr) {
         kernel::spmv<num_thread_per_worker, atomic>
diff --git a/hip/matrix/ell_kernels.hip.cpp b/hip/matrix/ell_kernels.hip.cpp
index 24966a64113..f1972635385 100644
--- a/hip/matrix/ell_kernels.hip.cpp
+++ b/hip/matrix/ell_kernels.hip.cpp
@@ -155,8 +155,8 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
         std::array<size_type, 1>{{num_stored_elements_per_row * stride}},
         a->get_const_values());
     const auto b_vals = gko::acc::range<b_accessor>(
-        std::array<size_type, 2>{{nrows, b->get_stride()}},
-        b->get_const_values());
+        std::array<size_type, 2>{{b->get_size()[0], b->get_size()[1]}},
+        b->get_const_values(), std::array<size_type, 1>{{b->get_stride()}});
 
     if (alpha == nullptr && beta == nullptr) {
         hipLaunchKernelGGL(