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 6ab824a
Show file tree
Hide file tree
Showing 4 changed files with 32 additions and 29 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
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
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 6ab824a

Please sign in to comment.