Expose ILU0 numeric boost configuration for CUDA and HIP by kakeueda · Pull Request #466 · ORNL/ReSolve

kakeueda · 2026-08-06T14:11:40Z

Description

The numeric boost tolerance and value in the CUDA ILU0 solver are hard-coded to 1e-6. This PR makes them configurable through the existing setCliParam() interface.

To see the effect of these values I tested the solver with FIDAP012 from the NIST Matrix Market, a nonsymmetric linear system arising from a fluid dynamics problem

./build/examples/sysGmres.exe \
  -m ~/tmp/matrixmarket/fidap012.mtx \
  -r ~/tmp/matrixmarket/fidap012_rhs1.mtx \
  -b cuda \
  -i fgmres

The number of iterations changed as follows:

Boost value	Number of FGMRES iterations	Final relative residual
`1e-6`	1121	`9.9392e-13`
`1e-7`	143	`9.2850e-13`
`1e-8`	143	`9.2850e-13`

Some problems are sensitive to this value.

Proposed changes

Added a zero_diagonal parameter to the CPU ILU0 solver through the existing setCliParam() interface.
Added boost_tolerance and boost_value parameters to the CUDA/HIP ILU0 solvers through the existing setCliParam() interface.
Added SystemSolver::getPreconditionerSolver() so users can configure ILU0 numerical stabilization parameters through SystemSolver.
Updated testSysGmres and MatrixFactorizationTests to configure ILU0 parameters.
Fixed misplaced entries in CHANGELOG.md.

Checklist

kakeueda · 2026-08-06T14:18:20Z

I’d appreciate any help with HIP testing

tamar-dewilde

Tested on HIP. All 64 tests and install validation passed.

tamar-dewilde · 2026-08-06T16:00:09Z


    LinSolverDirect&    getFactorizationSolver();
    LinSolverDirect&    getRefactorizationSolver();
+    LinSolverDirect&    getPreconditionerSolver();


The PR description says this getter lets users configure the numeric boost through SystemSolver. How should users call setZeroDiagonal() through it? The getter returns LinSolverDirect&, which does not declare that method, so the direct call does not compile.

I could only call the setter by including LinSolverDirectCuSparseILU0.hpp and casting the returned solver to that concrete type. Is that the intended use, or should LinSolverDirect provide a common way to access this setting? LinSolverDirectCpuILU0 already has the same setter.

Thank you for pointing this out! I added the BOOST_TOLERANCE and BOOST_VALUE parameter IDs and the corresponding cli param settings so users can now configure the numeric boost through SystemSolver rather than casting it to a concrete type.

I also implemented numeric boost support for HIP since it was needed for the configuration test added to testSysGmres.cpp.

Can you test this again on HIP?

kakeueda · 2026-08-07T09:20:56Z

One remaining inconsistency is that CPU ILU0 still uses zero_diagonal for both the zero-pivot threshold and replacement value. Separating them also requires changing CPU ILU0’s numerical behavior. I left it out of this PR.

shakedregev · 2026-08-07T14:26:57Z

You showed what the iteration number change to, what did it change from?

shakedregev

Looks good, neeed to test.

kakeueda · 2026-08-07T14:42:44Z

@shakedregev With the existing hard-coded settings (boost_tolerance = boost_value = 1e-6), the iteration count was 1121. Setting both parameters to either 1e-7 or 1e-8 reduced it to 143 iterations.

My Navier–Stokes FEM system also converges poorly with the default 1e-6, I suspect this value may be particularly problematic for some incompressible-flow systems (the system shown above comes from the Stokes equations as well).

shakedregev · 2026-08-07T14:56:33Z

Oh, now I understand.

shakedregev · 2026-08-07T15:12:17Z

Funnily my HIP test worked, but CUDA did not.
(The normal tests pass on both platforms).

./examples/sysGmres.exe   -m ../matrixmarket/fidap012.mtx   -r ../matrixmarket/fidap012_rhs1.mtx   -b cuda -t 1e-6
Matrix file: ../matrixmarket/fidap012.mtx
RHS file: ../matrixmarket/fidap012_rhs1.mtx
sysGmres with CUDA backend
solver.setMatrix returned status: 0
solver.preconditionerSetup returned status: 0
[ERROR] CUDA error in function cudaMemcpy(dst, src, sizeof(T) * n, cudaMemcpyDeviceToHost) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:133, error# 716
[ERROR] Dot product failed with error code 15
[ERROR] CUDA error in function cudaMemcpy(dst, src, sizeof(T) * n, cudaMemcpyHostToDevice) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:167, error# 716
[ERROR] scal returned error code 13
[ERROR] CUDA error in function cudaMemcpy(dst, src, sizeof(T) * n, cudaMemcpyDeviceToDevice) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:150, error# 716
[ERROR] Dot product failed with error code 15
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
solver.solve returned status: 0
	 Initial relative residual norm ||b-A*x||/||b|| : 1.0000000000000000e+00
	 Final relative residual norm   ||b-A*x||/||b|| : 0.0000000000000000e+00
	 Number of iterations                           : 1
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716
[ERROR] CUDA error in function cudaFree(v) at /home/szb/ReSolve_dir/ReSolve/resolve/cuda/CudaMemory.hpp:43, error# 716

tamar-dewilde · 2026-08-07T15:21:46Z


+    status_rocsparse_ = rocsparse_dcsrilu0_numeric_boost(workspace_->getRocsparseHandle(),
+                                                         info_A_,
+                                                         1,


Was it intentional to enable numeric boost by default on HIP here? I set this flag to 0 and reran the full HIP suite on Frontier, and all 64 tests still passed.

I enabled it for consistency with cuda but I agree that we can leave the hip settings as they are since we haven’t encountered the same issue there. Still the same issue could surface on hip as we expand testing to a broader range of systems.

A better approach would be to make numeric boost configurable on both backends (disabled by default on hip and enabled on cuda) so users can turn regularization on or off as needed.

kakeueda · 2026-08-07T21:13:59Z

@shakedregev Interesting, I couldn’t reproduce that error on my side (RTX 5050 Laptop GPU, CUDA Toolkit 12.9, driver 595.79, GCC 13). Could you try the same command on the develop branch and let me know whether it still occurs?

Just a quick note -t is not currently a valid option for sysGmres, unless you added it locally for convenience. The relevant parameters are not yet consistent between the cpu (zero_diagonal) and device (boost_tolerance and boost_value), so this setting is not currently exposed as a command-line option.

Running the same command but without -t, I got

ku@kakeru:~/dev/ReSolve$ ./build/examples/sysGmres.exe   -m ~/tmp/matrixmarket/fidap012.mtx   -r ~/tmp/matrixmarket/fidap012_rhs1.mtx -b cuda   
Matrix file: /home/ku/tmp/matrixmarket/fidap012.mtx
RHS file: /home/ku/tmp/matrixmarket/fidap012_rhs1.mtx
sysGmres with CUDA backend
solver.setMatrix returned status: 0
solver.preconditionerSetup returned status: 0
solver.solve returned status: 0
         Initial relative residual norm ||b-A*x||/||b|| : 1.0000000000000000e+00
         Final relative residual norm   ||b-A*x||/||b|| : 1.0808156599918938e-12
         Number of iterations                           : 966

kakeueda added 2 commits August 6, 2026 22:18

Expose CUDA ILU0 boost configuration

e99896d

Update CHANGELOG.md

0ca8727

kakeueda requested review from shakedregev and tamar-dewilde August 6, 2026 14:19

tamar-dewilde reviewed Aug 6, 2026

View reviewed changes

kakeueda added 2 commits August 7, 2026 08:59

Add CLI parameter support

124bee0

Split GPU ILU0 numeric boost configuration

ca8ea4a

kakeueda changed the title ~~Expose CUDA ILU0 boost configuration~~ Expose ILU0 numeric boost configuration for CUDA and HIP Aug 7, 2026

Rename numeric boost parameters

d9bf109

kakeueda requested a review from tamar-dewilde August 7, 2026 12:28

shakedregev reviewed Aug 7, 2026

View reviewed changes

shakedregev requested changes Aug 7, 2026

View reviewed changes

tamar-dewilde reviewed Aug 7, 2026

View reviewed changes

Navigation Menu

Uh oh!

Expose ILU0 numeric boost configuration for CUDA and HIP - #466

Expose ILU0 numeric boost configuration for CUDA and HIP#466
kakeueda wants to merge 5 commits into
developfrom
kakeru/add-setter-boost

kakeueda commented Aug 6, 2026 •

edited

Loading

Uh oh!

kakeueda commented Aug 6, 2026

Uh oh!

tamar-dewilde left a comment

Uh oh!

tamar-dewilde Aug 6, 2026

Uh oh!

kakeueda Aug 7, 2026 •

edited

Loading

Uh oh!

kakeueda commented Aug 7, 2026 •

edited

Loading

Uh oh!

shakedregev commented Aug 7, 2026

Uh oh!

shakedregev left a comment

Uh oh!

kakeueda commented Aug 7, 2026

Uh oh!

shakedregev commented Aug 7, 2026

Uh oh!

shakedregev commented Aug 7, 2026 •

edited

Loading

Uh oh!

tamar-dewilde Aug 7, 2026

Uh oh!

kakeueda Aug 7, 2026 •

edited

Loading

Uh oh!

kakeueda commented Aug 7, 2026 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

kakeueda commented Aug 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Proposed changes

Checklist

Uh oh!

kakeueda commented Aug 6, 2026

Uh oh!

tamar-dewilde left a comment

Choose a reason for hiding this comment

Uh oh!

tamar-dewilde Aug 6, 2026

Choose a reason for hiding this comment

Uh oh!

kakeueda Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kakeueda commented Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

shakedregev commented Aug 7, 2026

Uh oh!

shakedregev left a comment

Choose a reason for hiding this comment

Uh oh!

kakeueda commented Aug 7, 2026

Uh oh!

shakedregev commented Aug 7, 2026

Uh oh!

shakedregev commented Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

tamar-dewilde Aug 7, 2026

Choose a reason for hiding this comment

Uh oh!

kakeueda Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

kakeueda commented Aug 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

kakeueda commented Aug 6, 2026 •

edited

Loading

kakeueda Aug 7, 2026 •

edited

Loading

kakeueda commented Aug 7, 2026 •

edited

Loading

shakedregev commented Aug 7, 2026 •

edited

Loading

kakeueda Aug 7, 2026 •

edited

Loading

kakeueda commented Aug 7, 2026 •

edited

Loading