From 01bd3c4c4caae542c81c7abf58089072b291ffc9 Mon Sep 17 00:00:00 2001
From: Fritz Goebel <goebel.fritz@googlemail.com>
Date: Mon, 15 Mar 2021 13:43:42 +0100
Subject: [PATCH] use size_type for accessor dimensions

---
 cuda/matrix/ell_kernels.cu       | 7 +++----
 hip/matrix/ell_kernels.hip.cpp   | 7 +++----
 omp/matrix/ell_kernels.cpp       | 8 ++++----
 reference/matrix/ell_kernels.cpp | 8 ++++----
 4 files changed, 14 insertions(+), 16 deletions(-)

diff --git a/cuda/matrix/ell_kernels.cu b/cuda/matrix/ell_kernels.cu
index 33db65c4c1c..2dc669ba5a7 100644
--- a/cuda/matrix/ell_kernels.cu
+++ b/cuda/matrix/ell_kernels.cu
@@ -148,11 +148,10 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
                          b->get_size()[1], 1);
 
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{
-            {num_stored_elements_per_row * stride}},
+        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<long unsigned int, 2>{{nrows, b->get_stride()}},
+        std::array<size_type, 2>{{nrows, b->get_stride()}},
         b->get_const_values());
 
     if (alpha == nullptr && beta == nullptr) {
@@ -164,7 +163,7 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
                 as_cuda_type(c->get_values()), c->get_stride());
     } else if (alpha != nullptr && beta != nullptr) {
         const auto alpha_val = gko::acc::range<a_accessor>(
-            std::array<long unsigned int, 1>{1}, alpha->get_const_values());
+            std::array<size_type, 1>{1}, alpha->get_const_values());
         kernel::spmv<num_thread_per_worker, atomic>
             <<<grid_size, block_size, 0, 0>>>(
                 nrows, num_worker_per_row, as_cuda_accessor(alpha_val),
diff --git a/hip/matrix/ell_kernels.hip.cpp b/hip/matrix/ell_kernels.hip.cpp
index 8bcadd02370..24966a64113 100644
--- a/hip/matrix/ell_kernels.hip.cpp
+++ b/hip/matrix/ell_kernels.hip.cpp
@@ -152,11 +152,10 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
                          b->get_size()[1], 1);
 
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{
-            {num_stored_elements_per_row * stride}},
+        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<long unsigned int, 2>{{nrows, b->get_stride()}},
+        std::array<size_type, 2>{{nrows, b->get_stride()}},
         b->get_const_values());
 
     if (alpha == nullptr && beta == nullptr) {
@@ -168,7 +167,7 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
             b->get_stride(), as_hip_type(c->get_values()), c->get_stride());
     } else if (alpha != nullptr && beta != nullptr) {
         const auto alpha_val = gko::acc::range<a_accessor>(
-            std::array<long unsigned int, 1>{1}, alpha->get_const_values());
+            std::array<size_type, 1>{1}, alpha->get_const_values());
         hipLaunchKernelGGL(
             HIP_KERNEL_NAME(kernel::spmv<num_thread_per_worker, atomic>),
             dim3(grid_size), dim3(block_size), 0, 0, nrows, num_worker_per_row,
diff --git a/omp/matrix/ell_kernels.cpp b/omp/matrix/ell_kernels.cpp
index 76e7eea80ed..040176053d6 100644
--- a/omp/matrix/ell_kernels.cpp
+++ b/omp/matrix/ell_kernels.cpp
@@ -74,10 +74,10 @@ void spmv(std::shared_ptr<const OmpExecutor> exec,
     const auto stride = a->get_stride();
     const auto num_rows = b->get_size()[0];
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{num_stored_elements_per_row * stride},
+        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<long unsigned int, 2>{num_rows, b->get_stride()},
+        std::array<size_type, 2>{num_rows, b->get_stride()},
         b->get_const_values());
 
 #pragma omp parallel for
@@ -118,10 +118,10 @@ void advanced_spmv(std::shared_ptr<const OmpExecutor> exec,
     const auto stride = a->get_stride();
     const auto num_rows = b->get_size()[0];
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{num_stored_elements_per_row * stride},
+        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<long unsigned int, 2>{num_rows, b->get_stride()},
+        std::array<size_type, 2>{num_rows, b->get_stride()},
         b->get_const_values());
     const auto alpha_val = OutputValueType(alpha->at(0, 0));
     const auto beta_val = beta->at(0, 0);
diff --git a/reference/matrix/ell_kernels.cpp b/reference/matrix/ell_kernels.cpp
index ba6d687d18a..8b0f541799a 100644
--- a/reference/matrix/ell_kernels.cpp
+++ b/reference/matrix/ell_kernels.cpp
@@ -70,10 +70,10 @@ void spmv(std::shared_ptr<const ReferenceExecutor> exec,
     const auto stride = a->get_stride();
     const auto num_rows = b->get_size()[0];
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{num_stored_elements_per_row * stride},
+        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<long unsigned int, 2>{num_rows, b->get_stride()},
+        std::array<size_type, 2>{num_rows, b->get_stride()},
         b->get_const_values());
 
     for (size_type row = 0; row < a->get_size()[0]; row++) {
@@ -113,10 +113,10 @@ void advanced_spmv(std::shared_ptr<const ReferenceExecutor> exec,
     const auto stride = a->get_stride();
     const auto num_rows = b->get_size()[0];
     const auto a_vals = gko::acc::range<a_accessor>(
-        std::array<long unsigned int, 1>{num_stored_elements_per_row * stride},
+        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<long unsigned int, 2>{num_rows, b->get_stride()},
+        std::array<size_type, 2>{num_rows, b->get_stride()},
         b->get_const_values());
     const auto alpha_val = OutputValueType(alpha->at(0, 0));
     const auto beta_val = beta->at(0, 0);