Use thrust to generate the coarse of amgx_pgm #980

yhmtsai · 2022-03-04T10:37:50Z

This pull request eliminate the usage of spgemm in coarse grid generation of amgx_pgm
After pull request, the transposed is only needed by the weight matrix computation.
If the matrix is hermitian, the transpose can also be ignored.

upsj

LGTM! Minor nits. Also, this turns the previously deterministic generation (depending on SpGEMM determinism) into a non-deterministic one

upsj · 2022-03-04T10:43:04Z

cuda/multigrid/amgx_pgm_kernels.cu

+ thrust::device_pointer_cast(coarse_coo->get_row_idxs()),
+ thrust::device_pointer_cast(coarse_coo->get_col_idxs())));
+
+ thrust::reduce_by_key(thrust::device, key_it, key_it + fine_nnz, vals_it,


note: reduce_by_key can be non-deterministic

upsj · 2022-03-04T10:43:11Z

cuda/multigrid/amgx_pgm_kernels.cu

+ auto it = thrust::make_zip_iterator(
+ thrust::make_tuple(thrust::device_pointer_cast(row_idxs),
+ thrust::device_pointer_cast(col_idxs)));
+ thrust::stable_sort_by_key(thrust::device, it, it + nnz, vals_it);


nit: does this need to be stable?

Suggested change

thrust::stable_sort_by_key(thrust::device, it, it + nnz, vals_it);

thrust::sort_by_key(thrust::device, it, it + nnz, vals_it);

Yes only when reduce_by_key is determinstic. If reduce_by_key is not determinstic, this stable version does not help the coarse determinstic property

upsj · 2022-03-04T10:43:47Z

dpcpp/multigrid/amgx_pgm_kernels.dp.cpp

+ auto policy =
+ oneapi::dpl::execution::make_device_policy(*exec->get_queue());
+ auto it = oneapi::dpl::make_zip_iterator(row_idxs, col_idxs, vals);
+ std::stable_sort(policy, it, it + nnz, [](auto a, auto b) {


same here

Suggested change

std::stable_sort(policy, it, it + nnz, [](auto a, auto b) {

std::sort(policy, it, it + nnz, [](auto a, auto b) {

upsj · 2022-03-04T10:44:49Z

hip/multigrid/amgx_pgm_kernels.hip.cpp

+
+
+template <typename IndexType>
+void sort_agg(std::shared_ptr<const DefaultExecutor> exec, IndexType num,


can be pulled out into common/cuda_hip

upsj · 2022-03-04T10:45:07Z

omp/multigrid/amgx_pgm_kernels.cpp

+ IndexType* row_idxs, IndexType* col_idxs, ValueType* vals)
+{
+ auto it = detail::make_zip_iterator(row_idxs, col_idxs, vals);
+ std::stable_sort(it, it + nnz, [](auto a, auto b) {


Suggested change

std::stable_sort(it, it + nnz, [](auto a, auto b) {

std::sort(it, it + nnz, [](auto a, auto b) {

test/matrix/matrix.cpp

pratikvn

Really nice work! Do you have some performance numbers as to how much this improves the generation time over the previous approach ?

core/multigrid/amgx_pgm.cpp

hip/multigrid/amgx_pgm_kernels.hip.cpp

omp/multigrid/amgx_pgm_kernels.cpp

reference/multigrid/amgx_pgm_kernels.cpp

codecov · 2022-04-13T22:47:49Z

Codecov Report

Merging #980 (791f703) into develop (18c2a6a) will decrease coverage by 0.00%.
The diff coverage is 92.96%.

@@             Coverage Diff             @@
##           develop     #980      +/-   ##
===========================================
- Coverage    93.49%   93.49%   -0.01%     
===========================================
  Files          485      486       +1     
  Lines        40686    40797     +111     
===========================================
+ Hits         38041    38143     +102     
- Misses        2645     2654       +9

Impacted Files	Coverage Δ
core/device_hooks/common_kernels.inc.cpp	`0.00% <0.00%> (ø)`
test/base/device_matrix_data_kernels.cpp	`100.00% <ø> (ø)`
test/matrix/matrix.cpp	`96.01% <ø> (ø)`
common/unified/multigrid/amgx_pgm_kernels.cpp	`86.48% <90.90%> (+1.87%)`	⬆️
omp/multigrid/amgx_pgm_kernels.cpp	`96.66% <96.66%> (ø)`
reference/multigrid/amgx_pgm_kernels.cpp	`98.49% <97.87%> (-0.35%)`	⬇️
core/base/iterator_factory.hpp	`100.00% <100.00%> (ø)`
core/multigrid/amgx_pgm.cpp	`100.00% <100.00%> (ø)`

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 5b00508...791f703. Read the comment docs.

yhmtsai · 2022-04-14T22:27:04Z

rebase!

Co-authored-by: Pratik Nayak <[email protected]> Co-authored-by: Tobias Ribizel <[email protected]>

ginkgo-bot · 2022-04-14T22:38:59Z

Note: This PR changes the Ginkgo ABI:

Functions changes summary: 8 Removed, 0 Changed, 32 Added functions
Variables changes summary: 0 Removed, 0 Changed, 0 Added variable

For details check the full ABI diff under Artifacts here

yhmtsai · 2022-04-15T09:44:11Z

@pratikvn there are some result. The performance is not so good, and the generation should be improved more in the future. This version is mainly to avoid the SpGEMM calls such that we can implement it in distributed version more easily.
L-shape
original: 82 80 80
updated: 65 65 64
beam
original: 80 80 74
updated: 74 72 72

pratikvn · 2022-04-15T10:24:10Z

@yhmtsai, thanks for checking. I think that is still a significant improvement.

sonarcloud · 2022-04-15T11:53:25Z

SonarCloud Quality Gate failed.

0 Bugs
0 Vulnerabilities
0 Security Hotspots
29 Code Smells

72.1% Coverage
13.1% Duplication

Advertise release 1.5.0 and last changes + Add changelog, + Update third party libraries + A small fix to a CMake file See PR: #1195 The Ginkgo team is proud to announce the new Ginkgo minor release 1.5.0. This release brings many important new features such as: - MPI-based multi-node support for all matrix formats and most solvers; - full DPC++/SYCL support, - functionality and interface for GPU-resident sparse direct solvers, - an interface for wrapping solvers with scaling and reordering applied, - a new algebraic Multigrid solver/preconditioner, - improved mixed-precision support, - support for device matrix assembly, and much more. If you face an issue, please first check our [known issues page](https://github.com/ginkgo-project/ginkgo/wiki/Known-Issues) and the [open issues list](https://github.com/ginkgo-project/ginkgo/issues) and if you do not find a solution, feel free to [open a new issue](https://github.com/ginkgo-project/ginkgo/issues/new/choose) or ask a question using the [github discussions](https://github.com/ginkgo-project/ginkgo/discussions). Supported systems and requirements: + For all platforms, CMake 3.13+ + C++14 compliant compiler + Linux and macOS + GCC: 5.5+ + clang: 3.9+ + Intel compiler: 2018+ + Apple LLVM: 8.0+ + NVHPC: 22.7+ + Cray Compiler: 14.0.1+ + CUDA module: CUDA 9.2+ or NVHPC 22.7+ + HIP module: ROCm 4.0+ + DPC++ module: Intel OneAPI 2021.3 with oneMKL and oneDPL. Set the CXX compiler to `dpcpp`. + Windows + MinGW and Cygwin: GCC 5.5+ + Microsoft Visual Studio: VS 2019 + CUDA module: CUDA 9.2+, Microsoft Visual Studio + OpenMP module: MinGW or Cygwin. Algorithm and important feature additions: + Add MPI-based multi-node for all matrix formats and solvers (except GMRES and IDR). ([#676](#676), [#908](#908), [#909](#909), [#932](#932), [#951](#951), [#961](#961), [#971](#971), [#976](#976), [#985](#985), [#1007](#1007), [#1030](#1030), [#1054](#1054), [#1100](#1100), [#1148](#1148)) + Porting the remaining algorithms (preconditioners like ISAI, Jacobi, Multigrid, ParILU(T) and ParIC(T)) to DPC++/SYCL, update to SYCL 2020, and improve support and performance ([#896](#896), [#924](#924), [#928](#928), [#929](#929), [#933](#933), [#943](#943), [#960](#960), [#1057](#1057), [#1110](#1110), [#1142](#1142)) + Add a Sparse Direct interface supporting GPU-resident numerical LU factorization, symbolic Cholesky factorization, improved triangular solvers, and more ([#957](#957), [#1058](#1058), [#1072](#1072), [#1082](#1082)) + Add a ScaleReordered interface that can wrap solvers and automatically apply reorderings and scalings ([#1059](#1059)) + Add a Multigrid solver and improve the aggregation based PGM coarsening scheme ([#542](#542), [#913](#913), [#980](#980), [#982](#982), [#986](#986)) + Add infrastructure for unified, lambda-based, backend agnostic, kernels and utilize it for some simple kernels ([#833](#833), [#910](#910), [#926](#926)) + Merge different CUDA, HIP, DPC++ and OpenMP tests under a common interface ([#904](#904), [#973](#973), [#1044](#1044), [#1117](#1117)) + Add a device_matrix_data type for device-side matrix assembly ([#886](#886), [#963](#963), [#965](#965)) + Add support for mixed real/complex BLAS operations ([#864](#864)) + Add a FFT LinOp for all but DPC++/SYCL ([#701](#701)) + Add FBCSR support for NVIDIA and AMD GPUs and CPUs with OpenMP ([#775](#775)) + Add CSR scaling ([#848](#848)) + Add array::const_view and equivalent to create constant matrices from non-const data ([#890](#890)) + Add a RowGatherer LinOp supporting mixed precision to gather dense matrix rows ([#901](#901)) + Add mixed precision SparsityCsr SpMV support ([#970](#970)) + Allow creating CSR submatrix including from (possibly discontinuous) index sets ([#885](#885), [#964](#964)) + Add a scaled identity addition (M <- aI + bM) feature interface and impls for Csr and Dense ([#942](#942)) Deprecations and important changes: + Deprecate AmgxPgm in favor of the new Pgm name. ([#1149](#1149)). + Deprecate specialized residual norm classes in favor of a common `ResidualNorm` class ([#1101](#1101)) + Deprecate CamelCase non-polymorphic types in favor of snake_case versions (like array, machine_topology, uninitialized_array, index_set) ([#1031](#1031), [#1052](#1052)) + Bug fix: restrict gko::share to rvalue references (*possible interface break*) ([#1020](#1020)) + Bug fix: when using cuSPARSE's triangular solvers, specifying the factory parameter `num_rhs` is now required when solving for more than one right-hand side, otherwise an exception is thrown ([#1184](#1184)). + Drop official support for old CUDA < 9.2 ([#887](#887)) Improved performance additions: + Reuse tmp storage in reductions in solvers and add a mutable workspace to all solvers ([#1013](#1013), [#1028](#1028)) + Add HIP unsafe atomic option for AMD ([#1091](#1091)) + Prefer vendor implementations for Dense dot, conj_dot and norm2 when available ([#967](#967)). + Tuned OpenMP SellP, COO, and ELL SpMV kernels for a small number of RHS ([#809](#809)) Fixes: + Fix various compilation warnings ([#1076](#1076), [#1183](#1183), [#1189](#1189)) + Fix issues with hwloc-related tests ([#1074](#1074)) + Fix include headers for GCC 12 ([#1071](#1071)) + Fix for simple-solver-logging example ([#1066](#1066)) + Fix for potential memory leak in Logger ([#1056](#1056)) + Fix logging of mixin classes ([#1037](#1037)) + Improve value semantics for LinOp types, like moved-from state in cross-executor copy/clones ([#753](#753)) + Fix some matrix SpMV and conversion corner cases ([#905](#905), [#978](#978)) + Fix uninitialized data ([#958](#958)) + Fix CUDA version requirement for cusparseSpSM ([#953](#953)) + Fix several issues within bash-script ([#1016](#1016)) + Fixes for `NVHPC` compiler support ([#1194](#1194)) Other additions: + Simplify and properly name GMRES kernels ([#861](#861)) + Improve pkg-config support for non-CMake libraries ([#923](#923), [#1109](#1109)) + Improve gdb pretty printer ([#987](#987), [#1114](#1114)) + Add a logger highlighting inefficient allocation and copy patterns ([#1035](#1035)) + Improved and optimized test random matrix generation ([#954](#954), [#1032](#1032)) + Better CSR strategy defaults ([#969](#969)) + Add `move_from` to `PolymorphicObject` ([#997](#997)) + Remove unnecessary device_guard usage ([#956](#956)) + Improvements to the generic accessor for mixed-precision ([#727](#727)) + Add a naive lower triangular solver implementation for CUDA ([#764](#764)) + Add support for int64 indices from CUDA 11 onward with SpMV and SpGEMM ([#897](#897)) + Add a L1 norm implementation ([#900](#900)) + Add reduce_add for arrays ([#831](#831)) + Add utility to simplify Dense View creation from an existing Dense vector ([#1136](#1136)). + Add a custom transpose implementation for Fbcsr and Csr transpose for unsupported vendor types ([#1123](#1123)) + Make IDR random initilization deterministic ([#1116](#1116)) + Move the algorithm choice for triangular solvers from Csr::strategy_type to a factory parameter ([#1088](#1088)) + Update CUDA archCoresPerSM ([#1175](#1116)) + Add kernels for Csr sparsity pattern lookup ([#994](#994)) + Differentiate between structural and numerical zeros in Ell/Sellp ([#1027](#1027)) + Add a binary IO format for matrix data ([#984](#984)) + Add a tuple zip_iterator implementation ([#966](#966)) + Simplify kernel stubs and declarations ([#888](#888)) + Simplify GKO_REGISTER_OPERATION with lambdas ([#859](#859)) + Simplify copy to device in tests and examples ([#863](#863)) + More verbose output to array assertions ([#858](#858)) + Allow parallel compilation for Jacobi kernels ([#871](#871)) + Change clang-format pointer alignment to left ([#872](#872)) + Various improvements and fixes to the benchmarking framework ([#750](#750), [#759](#759), [#870](#870), [#911](#911), [#1033](#1033), [#1137](#1137)) + Various documentation improvements ([#892](#892), [#921](#921), [#950](#950), [#977](#977), [#1021](#1021), [#1068](#1068), [#1069](#1069), [#1080](#1080), [#1081](#1081), [#1108](#1108), [#1153](#1153), [#1154](#1154)) + Various CI improvements ([#868](#868), [#874](#874), [#884](#884), [#889](#889), [#899](#899), [#903](#903), [#922](#922), [#925](#925), [#930](#930), [#936](#936), [#937](#937), [#958](#958), [#882](#882), [#1011](#1011), [#1015](#1015), [#989](#989), [#1039](#1039), [#1042](#1042), [#1067](#1067), [#1073](#1073), [#1075](#1075), [#1083](#1083), [#1084](#1084), [#1085](#1085), [#1139](#1139), [#1178](#1178), [#1187](#1187))

yhmtsai self-assigned this Mar 4, 2022

ginkgo-bot added mod:all This touches all Ginkgo modules. reg:build This is related to the build system. type:matrix-format This is related to the Matrix formats type:multigrid This is related to multigrid labels Mar 4, 2022

upsj approved these changes Mar 4, 2022

View reviewed changes

yhmtsai mentioned this pull request Mar 4, 2022

Multigrid improvement and fix #982

Merged

2 tasks

yhmtsai added the 1:ST:ready-for-review This PR is ready for review label Mar 17, 2022

pratikvn approved these changes Mar 29, 2022

View reviewed changes

core/multigrid/amgx_pgm.cpp Outdated Show resolved Hide resolved

hip/multigrid/amgx_pgm_kernels.hip.cpp Outdated Show resolved Hide resolved

omp/multigrid/amgx_pgm_kernels.cpp Outdated Show resolved Hide resolved

reference/multigrid/amgx_pgm_kernels.cpp Show resolved Hide resolved

yhmtsai force-pushed the amgx_pgm_generation branch 2 times, most recently from b7c62ca to cd30f12 Compare April 13, 2022 13:31

yhmtsai added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels Apr 14, 2022

yhmtsai and others added 6 commits April 14, 2022 22:27

add agg to sparsitycsr

b7846b5

dpcpp: add agg to sparsitycsr

69e555d

generate the coarse csr without csr apply

ed10e3a

dpcpp: generate coarse csr

81bc24e

try adding the lambda

e25e00c

use sort when reduce isn't determinstic, rm unused, put into cuda_hip

791f703

Co-authored-by: Pratik Nayak <[email protected]> Co-authored-by: Tobias Ribizel <[email protected]>

ginkgo-bot force-pushed the amgx_pgm_generation branch from cd30f12 to 791f703 Compare April 14, 2022 22:27

yhmtsai merged commit b7a8edf into develop Apr 15, 2022

yhmtsai deleted the amgx_pgm_generation branch April 15, 2022 13:14

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Use thrust to generate the coarse of amgx_pgm #980

Use thrust to generate the coarse of amgx_pgm #980

yhmtsai commented Mar 4, 2022

upsj left a comment

upsj Mar 4, 2022

upsj Mar 4, 2022

yhmtsai Mar 7, 2022

upsj Mar 4, 2022

upsj Mar 4, 2022

upsj Mar 4, 2022

pratikvn left a comment

codecov bot commented Apr 13, 2022 •

edited

Loading

yhmtsai commented Apr 14, 2022

ginkgo-bot commented Apr 14, 2022

yhmtsai commented Apr 15, 2022

pratikvn commented Apr 15, 2022

sonarcloud bot commented Apr 15, 2022

	thrust::stable_sort_by_key(thrust::device, it, it + nnz, vals_it);
	thrust::sort_by_key(thrust::device, it, it + nnz, vals_it);

	std::stable_sort(policy, it, it + nnz, [](auto a, auto b) {
	std::sort(policy, it, it + nnz, [](auto a, auto b) {



		template <typename IndexType>
		void sort_agg(std::shared_ptr<const DefaultExecutor> exec, IndexType num,

	std::stable_sort(it, it + nnz, [](auto a, auto b) {
	std::sort(it, it + nnz, [](auto a, auto b) {

Use thrust to generate the coarse of amgx_pgm #980

Use thrust to generate the coarse of amgx_pgm #980

Conversation

yhmtsai commented Mar 4, 2022

upsj left a comment

Choose a reason for hiding this comment

upsj Mar 4, 2022

Choose a reason for hiding this comment

upsj Mar 4, 2022

Choose a reason for hiding this comment

yhmtsai Mar 7, 2022

Choose a reason for hiding this comment

upsj Mar 4, 2022

Choose a reason for hiding this comment

upsj Mar 4, 2022

Choose a reason for hiding this comment

upsj Mar 4, 2022

Choose a reason for hiding this comment

pratikvn left a comment

Choose a reason for hiding this comment

codecov bot commented Apr 13, 2022 • edited Loading

Codecov Report

yhmtsai commented Apr 14, 2022

ginkgo-bot commented Apr 14, 2022

yhmtsai commented Apr 15, 2022

pratikvn commented Apr 15, 2022

sonarcloud bot commented Apr 15, 2022

codecov bot commented Apr 13, 2022 •

edited

Loading