Skip to content

Commit

Permalink
Code Review
Browse files Browse the repository at this point in the history
  • Loading branch information
@fritzgoebel
fritzgoebel committed Mar 18, 2021
1 parent 7a9baa5 commit ba95673
Show file tree
Hide file tree
Showing 6 changed files with 38 additions and 36 deletions.
18 changes: 9 additions & 9 deletions common/matrix/ell_kernels.hpp.inc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -41,8 +41,8 @@ __device__ void spmv_kernel(
const size_type num_rows, const int num_worker_per_row,
acc::range<a_accessor> val, const IndexType *__restrict__ col,
const size_type stride, const size_type num_stored_elements_per_row,
acc::range<b_accessor> b, const size_type b_stride,
OutputValueType *__restrict__ c, const size_type c_stride, Closure op)
acc::range<b_accessor> b, OutputValueType *__restrict__ c,
const size_type c_stride, Closure op)
{
const auto tidx = thread::get_thread_id_flat();
const decltype(tidx) column_id = blockIdx.y;
Expand Down Expand Up @@ -109,12 +109,12 @@ __global__ __launch_bounds__(default_block_size) void spmv(
const size_type num_rows, const int num_worker_per_row,
acc::range<a_accessor> val, const IndexType *__restrict__ col,
const size_type stride, const size_type num_stored_elements_per_row,
acc::range<b_accessor> b, const size_type b_stride,
OutputValueType *__restrict__ c, const size_type c_stride)
acc::range<b_accessor> b, OutputValueType *__restrict__ c,
const size_type c_stride)
{
spmv_kernel<num_thread_per_worker, atomic>(
num_rows, num_worker_per_row, val, col, stride,
num_stored_elements_per_row, b, b_stride, c, c_stride,
num_stored_elements_per_row, b, c, c_stride,
[](const OutputValueType &x, const OutputValueType &y) { return x; });
}

Expand All @@ -126,8 +126,8 @@ __global__ __launch_bounds__(default_block_size) void spmv(
acc::range<a_accessor> alpha, acc::range<a_accessor> val,
const IndexType *__restrict__ col, const size_type stride,
const size_type num_stored_elements_per_row, acc::range<b_accessor> b,
const size_type b_stride, const OutputValueType *__restrict__ beta,
OutputValueType *__restrict__ c, const size_type c_stride)
const OutputValueType *__restrict__ beta, OutputValueType *__restrict__ c,
const size_type c_stride)
{
const OutputValueType alpha_val = alpha(0);
const OutputValueType beta_val = beta[0];
Expand All @@ -138,14 +138,14 @@ __global__ __launch_bounds__(default_block_size) void spmv(
if (atomic) {
spmv_kernel<num_thread_per_worker, atomic>(
num_rows, num_worker_per_row, val, col, stride,
num_stored_elements_per_row, b, b_stride, c, c_stride,
num_stored_elements_per_row, b, c, c_stride,
[&alpha_val](const OutputValueType &x, const OutputValueType &y) {
return alpha_val * x;
});
} else {
spmv_kernel<num_thread_per_worker, atomic>(
num_rows, num_worker_per_row, val, col, stride,
num_stored_elements_per_row, b, b_stride, c, c_stride,
num_stored_elements_per_row, b, c, c_stride,
[&alpha_val, &beta_val](const OutputValueType &x,
const OutputValueType &y) {
return alpha_val * x + beta_val * y;
Expand Down
6 changes: 3 additions & 3 deletions cuda/matrix/ell_kernels.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -159,8 +159,8 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
<<<grid_size, block_size, 0, 0>>>(
nrows, num_worker_per_row, as_cuda_accessor(a_vals),
a->get_const_col_idxs(), stride, num_stored_elements_per_row,
as_cuda_accessor(b_vals), b->get_stride(),
as_cuda_type(c->get_values()), c->get_stride());
as_cuda_accessor(b_vals), 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<size_type, 1>{1}, alpha->get_const_values());
Expand All @@ -169,7 +169,7 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
nrows, num_worker_per_row, as_cuda_accessor(alpha_val),
as_cuda_accessor(a_vals), a->get_const_col_idxs(), stride,
num_stored_elements_per_row, as_cuda_accessor(b_vals),
b->get_stride(), as_cuda_type(beta->get_const_values()),
as_cuda_type(beta->get_const_values()),
as_cuda_type(c->get_values()), c->get_stride());
} else {
GKO_KERNEL_NOT_FOUND;
Expand Down
27 changes: 15 additions & 12 deletions dpcpp/matrix/ell_kernels.dp.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,24 +56,27 @@ namespace dpcpp {
namespace ell {


template <typename ValueType, typename IndexType>
template <typename InputValueType, typename MatrixValueType,
typename OutputValueType, typename IndexType>
void spmv(std::shared_ptr<const DpcppExecutor> exec,
const matrix::Ell<ValueType, IndexType> *a,
const matrix::Dense<ValueType> *b,
matrix::Dense<ValueType> *c) GKO_NOT_IMPLEMENTED;
const matrix::Ell<MatrixValueType, IndexType> *a,
const matrix::Dense<InputValueType> *b,
matrix::Dense<OutputValueType> *c) GKO_NOT_IMPLEMENTED;

GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_ELL_SPMV_KERNEL);
GKO_INSTANTIATE_FOR_EACH_MIXED_VALUE_AND_INDEX_TYPE(
GKO_DECLARE_ELL_SPMV_KERNEL);


template <typename ValueType, typename IndexType>
template <typename InputValueType, typename MatrixValueType,
typename OutputValueType, typename IndexType>
void advanced_spmv(std::shared_ptr<const DpcppExecutor> exec,
const matrix::Dense<ValueType> *alpha,
const matrix::Ell<ValueType, IndexType> *a,
const matrix::Dense<ValueType> *b,
const matrix::Dense<ValueType> *beta,
matrix::Dense<ValueType> *c) GKO_NOT_IMPLEMENTED;
const matrix::Dense<MatrixValueType> *alpha,
const matrix::Ell<MatrixValueType, IndexType> *a,
const matrix::Dense<InputValueType> *b,
const matrix::Dense<OutputValueType> *beta,
matrix::Dense<OutputValueType> *c) GKO_NOT_IMPLEMENTED;

GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
GKO_INSTANTIATE_FOR_EACH_MIXED_VALUE_AND_INDEX_TYPE(
GKO_DECLARE_ELL_ADVANCED_SPMV_KERNEL);


Expand Down
7 changes: 3 additions & 4 deletions hip/matrix/ell_kernels.hip.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -164,7 +164,7 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
dim3(grid_size), dim3(block_size), 0, 0, nrows, num_worker_per_row,
as_hip_accessor(a_vals), a->get_const_col_idxs(), stride,
num_stored_elements_per_row, as_hip_accessor(b_vals),
b->get_stride(), as_hip_type(c->get_values()), c->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<size_type, 1>{1}, alpha->get_const_values());
Expand All @@ -173,9 +173,8 @@ void abstract_spmv(syn::value_list<int, info>, int num_worker_per_row,
dim3(grid_size), dim3(block_size), 0, 0, nrows, num_worker_per_row,
as_hip_accessor(alpha_val), as_hip_accessor(a_vals),
a->get_const_col_idxs(), stride, num_stored_elements_per_row,
as_hip_accessor(b_vals), b->get_stride(),
as_hip_type(beta->get_const_values()), as_hip_type(c->get_values()),
c->get_stride());
as_hip_accessor(b_vals), as_hip_type(beta->get_const_values()),
as_hip_type(c->get_values()), c->get_stride());
} else {
GKO_KERNEL_NOT_FOUND;
}
Expand Down
8 changes: 4 additions & 4 deletions omp/matrix/ell_kernels.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -77,8 +77,8 @@ void spmv(std::shared_ptr<const OmpExecutor> exec,
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>{num_rows, 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()}});

#pragma omp parallel for
for (size_type row = 0; row < a->get_size()[0]; row++) {
Expand Down Expand Up @@ -121,8 +121,8 @@ void advanced_spmv(std::shared_ptr<const OmpExecutor> exec,
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>{num_rows, 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()}});
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
View file
Original file line number Diff line number Diff line change
Expand Up @@ -73,8 +73,8 @@ void spmv(std::shared_ptr<const ReferenceExecutor> exec,
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>{num_rows, 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()}});

for (size_type row = 0; row < a->get_size()[0]; row++) {
for (size_type j = 0; j < c->get_size()[1]; j++) {
Expand Down Expand Up @@ -116,8 +116,8 @@ void advanced_spmv(std::shared_ptr<const ReferenceExecutor> exec,
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>{num_rows, 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()}});
const auto alpha_val = OutputValueType(alpha->at(0, 0));
const auto beta_val = beta->at(0, 0);

Expand Down

0 comments on commit ba95673

Please sign in to comment.