cuda_kernels_mcsr.hpp

Go to the documentation of this file.

#ifndef PARALUTION_GPU_CUDA_KERNELS_MCSR_HPP_
#define PARALUTION_GPU_CUDA_KERNELS_MCSR_HPP_

#include "../matrix_formats_ind.hpp"

namespace paralution {

template <typename ValueType, typename IndexType>
__global__ void kernel_mcsr_spmv_scalar(const IndexType nrow, const IndexType *row_offset, 
                                        const IndexType *col, const ValueType *val, 
                                        const ValueType *in, ValueType *out) {

  IndexType ai = blockIdx.x*blockDim.x+threadIdx.x;
  IndexType aj;

  if (ai <nrow) {

    ValueType sum = val[ai] * in[ai];

    for (aj=row_offset[ai]; aj<row_offset[ai+1]; ++aj)
      sum = sum + val[aj]*in[col[aj]];

    out[ai] = sum;

  }

}


template <typename ValueType, typename IndexType>
__global__ void kernel_mcsr_add_spmv_scalar(const IndexType nrow, const IndexType *row_offset, 
                                            const IndexType *col, const ValueType *val, 
                                            const ValueType scalar,
                                            const ValueType *in, ValueType *out) {

  IndexType ai = blockIdx.x*blockDim.x+threadIdx.x;
  IndexType aj;

  if (ai <nrow) {

    out[ai] = out[ai] + scalar*val[ai] * in[ai];

    for (aj=row_offset[ai]; aj<row_offset[ai+1]; ++aj) {
      out[ai] = out[ai] + scalar*val[aj]*in[col[aj]];
    }

  }
}


}

#endif