dpcpp ports amgx_pgm

[yhmtsai]

This PR ports the amgx_pgm to dpcpp.
the current thrust::tie can not be translated to dpcpp by dpct.

@upsj Do you know something like thrust::tie inside oneDPL?

[upsj]

std::tie from <tuple>

[Slaedr]

LGTM! Looks completely consistent with the cuda implementation. Just one comment below.

[Slaedr]

Just putting the kernel call directly here makes the code easier to navigate and reduces code size. Is it simply that DPCT creates the extra function by default?

Suggested change

	kernel::match_edge_kernel(grid, default_block_size, 0, exec->get_queue(),
	num, strongest_neighbor.get_const_data(),
	agg.get_data());
	exec->get_queue()->parallel_for(
	sycl_nd_range(grid, block), [&](sycl::nd_item<3> item_ct1) {
	kernel::match_edge_kernel(num, strongest_neighbor.get_const_data(),
	agg_vals.get_data(), item_ct1);
	});

[upsj]

LGTM as well! I agree with @Slaedr, and would even go one step further: All of these kernels use one thread per row. So you could simplify them heavily by putting them directly into a 1D parallel_for

[yhmtsai]

no, the extra is our own effort.
it gives the same structure as the cuda and hip call including the share memory usage.
when the kernel call many times, the function in dpcpp is hard to align with other executor

[codecov]

Codecov Report

Merging #933 (905f6c1) into develop (4a7e36f) will increase coverage by 0.00%.
The diff coverage is 92.85%.

❗ Current head 905f6c1 differs from pull request most recent head 1d16ba1. Consider uploading reports for the commit 1d16ba1 to get more accurate results

@@            Coverage Diff            @@
##           develop     #933    +/-   ##
=========================================
  Coverage    93.36%   93.37%            
=========================================
  Files          476      476            
  Lines        39368    39244   -124     
=========================================
- Hits         36755    36643   -112     
+ Misses        2613     2601    -12     

Impacted Files	Coverage Δ
omp/base/kernel_launch.hpp	88.46% <ø> (ø)
common/unified/multigrid/amgx_pgm_kernels.cpp	84.61% <84.61%> (ø)
test/multigrid/amgx_pgm_kernels.cpp	100.00% <100.00%> (ø)
include/ginkgo/core/base/matrix_data.hpp	97.35% <0.00%> (-0.23%)	⬇️
core/base/device_matrix_data.cpp	100.00% <0.00%> (ø)
omp/base/device_matrix_data_kernels.cpp	100.00% <0.00%> (ø)
core/device_hooks/common_kernels.inc.cpp	0.00% <0.00%> (ø)
test/base/device_matrix_data_kernels.cpp	100.00% <0.00%> (ø)
reference/base/device_matrix_data_kernels.cpp	100.00% <0.00%> (ø)
include/ginkgo/core/base/device_matrix_data.hpp	62.50% <0.00%> (ø)
... and 3 more

---

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 e1f17e4...1d16ba1. Read the comment docs.

[yhmtsai]

@upsj @Slaedr I move those kernels to unified and add device_std as thrust for cuda/hip and as std for others

[Slaedr]

Excellent work!

Did you try benchmarking your previous implementations against these ones? Since we seem to be discussing multigrid performance quite often, it would be good to document the performance difference between the separate kernels and the unified ones. I guess it's unlikely there's a big difference, but it will be good to document it.

It might make sense to develop more fine-grain parallel kernels for the aggregation kernels, primarily to get good coalescing access, though I understand that is far from trivial. That can be left to another PR.

[Slaedr]

Right now, this algorithm (a) does not give coalesced access to weight matrix values and (b) processes the nonzeros in each row sequentially. Is that fine? Would you look into a more fine-grain parallel algorithm at some point? I guess this kernel is an important one, though I see that it's difficult to parallelize this further; it looks like you need a creative reduction. Similar for the next kernel.

[upsj]

I tend to agree here - there is a clear path to using warp-parallelism here in the future. It is a cheap way to implement it in DPC++. There may be a few abstract primitives we could use for the GPU kernels: Csr row-wise reduction + counting, Csr row-wise filtering, ... Might make sense to put that into a file in components? But for now I wouldn't block the PR for this.
Probably should take a look at specializing it before the release?

[yhmtsai]

yes it should be parallelized by warp per row.
we do not yet test the performance of the generation step.
at least, the RAP also need to improve

[Slaedr]

My only concern is we might have to go back to backend-specific kernels if we need more advanced algorithms. But for now, I think we should move ahead with this.

[Slaedr]

All these tests can now be move to tests/multigrid, I guess? Then you could remove hip/test/multigrid/amgx_pgm_kernels.cpp etc.

[yhmtsai]

unified ~107.86ms

4.81%  3.1193ms         4  779.82us  129.02us  1.8408ms find_strongest_neighbor
0.05%  33.488us         4  8.3720us  6.4610us  10.141us match_edge

original ~107.108 (skip_sorting 62.4452)

4.82%  3.1674ms         4  791.86us  146.08us  1.8393ms find_strongest_neighbor
0.06%  36.687us         4  9.1710us  6.8120us  10.939us match_edge

only check mfem l-shape 1 level generation on V100
it does not give noticeable difference between these.
the main bottleneck is spgemm and sorting.
sorting takes ~ 40 ~ 45% of generation and spgemm takes ~10%
without sorting. spgemm (sum of RadixSort and spgemm_core) takes ~20% and find_strongest_neighbor takes ~10%

[upsj]

I guess this should be reverted?

Suggested change

	std::ranlux48 rand_engine;
	std::default_random_engine rand_engine;

[yhmtsai]

good catch!

[sonarcloud]

SonarCloud Quality Gate failed.   

0 Bugs
0 Vulnerabilities
0 Security Hotspots
65 Code Smells

91.8% Coverage
7.2% Duplication