Cuda and Omp matrix conversion of Csr and Hybrid and Omp count_nonzeros of Ell

[yhmtsai]

Cuda and Omp matrix conversion of Csr and Hybrid and Ell count_nonzeros on Omp

Two functions are implemented in header files, so they are still on the host.

1. make_srow of csr with load_balance algorithm
2. the function of hybrid which decide the column limit of ell (It needs a sorting kernel)

I would like to put them in another PR if they are needed on Cuda.
I am not sure whether performance of make_srow is better or not.
The performance of sorting on GPU is better than CPU according to this link, so I think we need it on GPU.

[tcojean]

LGTM in general, I only have comments for reducing code duplication in the CUDA kernels, and checking quickly the quality of the OpenMP parallelization:

1. is using OpenMP with 1 thread creating a significant slowdown compared to reference?
2. is using OpenMP with 4 threads (or more) creating a significant slowdown compared to 1 thread?
   If that is the case, some internal loop parallelization might be the problems. If there is no such problem, it's good as it is.

[thoasm]

Overall, it looks good, but please get rid of the unsigned long long int and remove the code duplication.

[yhmtsai]

@tcojean @thoasm Both of you mention the SizeType (cuda_size_type or unsigned long long int).
I would like to get the exact nnz per row of Coo by spmv which uses atomicAdd.
The answer should be IndexType (int32 or int64) or size_type.
However, the cuda default atomicAdd of integer operator is for int, unsigned int, unsigned long long int.
I guess cuda always thinks unsigned long long int for 64bit. (it is also used in double atomicAdd)
In my environment, the compiler thinks size_type and unsigned long long int are different
Thus, I use unsigned long long int to handle int32, int64 or size_type.
I think I will do the converter for type with 64bit <-> unsigned long long int and implement its atomicAdd, like double atomicAdd using __longlong_as_double and __double_as_longlong in some architecture.

the function converter is like

__device__ size_type ull_as_size_type(ull value) {
    union DataType {
        unsigned long long int ull_val;
        size_type st_val;
    };
    DataType temp;
    temp.ull_val = value;
    return temp.st_val;
}

Do you think it works? or any idea for it.

[thoasm]

@yhmtsai Yes, this should work. I did something similar multiple times already when splitting double in various parts.
As far as I know, this (or something very similar with reinterpret_cast) would also work:

__device__ size_type convert(unsigned long long int val) {
    return reinterpret_cast<size_type &>(val);
}

I haven't benchmarked these two against each other (especially without doing any splitting), but I think they should both have almost no overhead (I actually hope they have none).

[yhmtsai]

I check the performance of omp kernel.
They are not faster than reference kernel.
I use 10000 x 2019 with at least 30 nonzero for testing.
Should I delete them or use the reference kernel for conversion performance test?

[yhmtsai]

• format_conversion.hpp is for handling the same function in format_conversion.
  convert_coo_row_idxs_to_ptrs is not in coo_kernel.hpp, and it is not directly used by coo.cpp.
  I also put it in format_conversion.hpp now.
• atomic.cuh should handle 32-bit and 64-bit atomic add now.
  if using cuda 10.1 or higher, it can also handle 16-bit atomic add.
  if using cuda 10.0 or higher, I bind the cuda function of __half and __half according to architecture.

[codecov]

Codecov Report

Merging #310 into develop will increase coverage by 0.12%.
The diff coverage is 100%.

@@             Coverage Diff             @@
##           develop     #310      +/-   ##
===========================================
+ Coverage    98.05%   98.17%   +0.12%     
===========================================
  Files          216      215       -1     
  Lines        16328    16529     +201     
===========================================
+ Hits         16010    16227     +217     
+ Misses         318      302      -16

Impacted Files	Coverage Δ
include/ginkgo/core/matrix/hybrid.hpp	91.52% <100%> (ø)	⬆️
omp/matrix/ell_kernels.cpp	91.66% <100%> (+6.66%)	⬆️
omp/test/matrix/csr_kernels.cpp	100% <100%> (ø)	⬆️
cuda/test/matrix/csr_kernels.cpp	100% <100%> (ø)	⬆️
omp/matrix/csr_kernels.cpp	94.28% <100%> (+7.92%)	⬆️
omp/test/matrix/ell_kernels.cpp	100% <100%> (ø)	⬆️
cuda/test/matrix/hybrid_kernels.cpp	100% <100%> (ø)	⬆️
omp/matrix/hybrid_kernels.cpp	98.82% <100%> (+20.56%)	⬆️
omp/test/matrix/hybrid_kernels.cpp	100% <100%> (ø)	⬆️

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update de51ee9...c255b78. Read the comment docs.

[tcojean]

LGTM. Good job on removing all these duplications.

[tcojean]

@yhmtsai for the OpenMP kernels, if they are not much slower than the reference ones, then I think it is fine. It at least shows the idea for parallelizing this. If there is a very important slowdown, then it's better to remove them.

[yhmtsai]

Hybrid fills all coo values, so it does not need to initial coo array.
I delete the functions of cuda and omp.
The following is omp performance.
The matrices are 10000 x 2019 with at least 30 elements in each row.

Csr-> Hybrid

• OMP_NUM_THREADS=1 ./csr_kernels
  Ref : 1.36327e+08ns.
  Omp : 1.09301e+08ns.

• OMP_NUM_THREADS=4 ./csr_kernels
  Ref : 1.33657e+08ns.
  Omp : 8.19972e+07ns. (Speedup 1.33x over 1-thread)

• OMP_NUM_THREADS=8 ./csr_kernels
  Ref : 1.32982e+08ns.
  Omp : 7.30211e+07ns. (Speedup 1.5x over 1-thread)

• OMP_NUM_THREADS=16 ./csr_kernels
  Ref : 1.3298e+08ns.
  Omp : 1.42636e+08ns. (Speedup 0.77x over 1-thread)

Hybrid -> Csr

• OMP_NUM_THREADS=1 ./hybrid_kernels
  Ref : 6.49456e+07ns.
  Omp : 9.57758e+07ns.

• OMP_NUM_THREADS=4 ./hybrid_kernels
  Ref : 6.49526e+07ns.
  Omp : 1.00031e+08ns. (Speedup 0.96x over 1-thread)

• OMP_NUM_THREADS=8 ./hybrid_kernels
  Ref : 7.43198e+07ns.
  Omp : 1.18494e+08ns. (Speedup 0.81x over 1-thread)

• OMP_NUM_THREADS=16 ./hybrid_kernels
  Ref : 6.39154e+07ns.
  Omp : 6.36856e+08ns. (Speedup 0.15x over 1-thread)

[thoasm]

It looks a lot better now, but there are some parts which I would like to see improved.
Mainly renaming the as function and adding some documentation.

[thoasm]

I think after replacing the GKO_ASSERT in the atomic_add (I don't think we need it anymore after the fixed static_assert is in place) and adding a bit more documentation, we should be ready to merge.

Well done!

[thoasm]

Regarding the runtime results, I think they are perfectly fine. I don't think it is necessary to change it to just use the reference implementation. The only issue could be both 16 thread cases, but I think there, it was too much overhead compared to the (relatively small) size of the matrix.

[thoasm]

LGTM!