Skip to content

Commit

Permalink
use size_type for accessor dimensions
Browse files Browse the repository at this point in the history
  • Loading branch information
fritzgoebel committed Mar 18, 2021
1 parent dd4e807 commit c07a6bc
Show file tree
Hide file tree
Showing 4 changed files with 14 additions and 16 deletions.
7 changes: 3 additions & 4 deletions cuda/matrix/ell_kernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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),
Expand Down
7 changes: 3 additions & 4 deletions hip/matrix/ell_kernels.hip.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Expand All @@ -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,
Expand Down
8 changes: 4 additions & 4 deletions omp/matrix/ell_kernels.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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);
Expand Down
8 changes: 4 additions & 4 deletions reference/matrix/ell_kernels.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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++) {
Expand Down Expand Up @@ -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);
Expand Down

0 comments on commit c07a6bc

Please sign in to comment.