48 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
54 real(kind=
rp),
allocatable :: w(:)
55 real(kind=
rp),
allocatable :: r(:)
56 real(kind=
rp),
allocatable :: p(:)
57 real(kind=
rp),
allocatable :: z(:)
58 type(c_ptr) :: w_d = c_null_ptr
59 type(c_ptr) :: r_d = c_null_ptr
60 type(c_ptr) :: p_d = c_null_ptr
61 type(c_ptr) :: z_d = c_null_ptr
62 type(c_ptr) :: gs_event = c_null_ptr
75 class(
pc_t),
optional,
intent(in),
target :: M
76 integer,
intent(in) :: n
77 integer,
intent(in) :: max_iter
78 real(kind=
rp),
optional,
intent(in) :: rel_tol
79 real(kind=
rp),
optional,
intent(in) :: abs_tol
80 logical,
optional,
intent(in) :: monitor
98 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
99 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
100 else if (
present(rel_tol) .and.
present(abs_tol))
then
101 call this%ksp_init(max_iter, rel_tol, abs_tol)
102 else if (
present(monitor) .and.
present(abs_tol))
then
103 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
104 else if (
present(rel_tol) .and.
present(monitor))
then
105 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
106 else if (
present(rel_tol))
then
107 call this%ksp_init(max_iter, rel_tol = rel_tol)
108 else if (
present(abs_tol))
then
109 call this%ksp_init(max_iter, abs_tol = abs_tol)
110 else if (
present(monitor))
then
111 call this%ksp_init(max_iter, monitor = monitor)
113 call this%ksp_init(max_iter)
125 if (
allocated(this%w))
then
126 if (c_associated(this%w_d))
then
132 if (
allocated(this%r))
then
133 if (c_associated(this%r_d))
then
139 if (
allocated(this%p))
then
140 if (c_associated(this%p_d))
then
146 if (
allocated(this%z))
then
147 if (c_associated(this%z_d))
then
155 if (c_associated(this%gs_event))
then
165 class(
ax_t),
intent(in) :: ax
166 type(
field_t),
intent(inout) :: x
167 integer,
intent(in) :: n
168 real(kind=
rp),
dimension(n),
intent(in) :: f
169 type(
coef_t),
intent(inout) :: coef
171 type(
gs_t),
intent(inout) :: gs_h
173 integer,
optional,
intent(in) :: niter
174 real(kind=
rp),
parameter :: one = 1.0
175 real(kind=
rp),
parameter :: zero = 0.0
176 integer :: iter, max_iter
177 real(kind=
rp) :: rnorm, rtr, rtr0, rtz2, rtz1
178 real(kind=
rp) :: beta, pap, alpha, alphm, norm_fac
183 if (
present(niter))
then
186 max_iter = this%max_iter
188 norm_fac = one/sqrt(coef%volume)
196 rnorm = sqrt(rtr)*norm_fac
197 ksp_results%res_start = rnorm
198 ksp_results%res_final = rnorm
200 if (
abscmp(rnorm, zero))
then
201 ksp_results%converged = .true.
204 call this%monitor_start(
'CG')
205 do iter = 1, max_iter
206 call this%M%solve(this%z, this%r, n)
210 if (iter .eq. 1) beta = zero
213 call ax%compute(this%w, this%p, coef, x%msh, x%Xh)
214 call gs_h%op(this%w, n, gs_op_add, this%gs_event)
216 call blst%apply(this%w, n)
226 if (iter .eq. 1) rtr0 = rtr
227 rnorm = sqrt(rtr)*norm_fac
228 call this%monitor_iter(iter, rnorm)
229 if (rnorm .lt. this%abs_tol)
then
233 call this%monitor_stop()
234 ksp_results%res_final = rnorm
235 ksp_results%iter = iter
236 ksp_results%converged = this%is_converged(iter, rnorm)
242 n, coef, blstx, blsty, blstz, gs_h, niter)
result(ksp_results)
244 class(
ax_t),
intent(in) :: ax
245 type(
field_t),
intent(inout) :: x
246 type(
field_t),
intent(inout) :: y
247 type(
field_t),
intent(inout) :: z
248 integer,
intent(in) :: n
249 real(kind=
rp),
dimension(n),
intent(in) :: fx
250 real(kind=
rp),
dimension(n),
intent(in) :: fy
251 real(kind=
rp),
dimension(n),
intent(in) :: fz
252 type(
coef_t),
intent(inout) :: coef
256 type(
gs_t),
intent(inout) :: gs_h
258 integer,
optional,
intent(in) :: niter
260 ksp_results(1) = this%solve(ax, x, fx, n, coef, blstx, gs_h, niter)
261 ksp_results(2) = this%solve(ax, y, fy, n, coef, blsty, gs_h, niter)
262 ksp_results(3) = this%solve(ax, z, fz, n, coef, blstz, gs_h, niter)
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
Return the device pointer for an associated Fortran array.
Map a Fortran array to a device (allocate and associate)
Unmap a Fortran array from a device (deassociate and free)
Defines a Matrix-vector product.
Defines various Conjugate Gradient methods for accelerators.
type(ksp_monitor_t) function cg_device_solve(this, ax, x, f, n, coef, blst, gs_h, niter)
Standard PCG solve.
subroutine cg_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device based PCG solver.
subroutine cg_device_free(this)
Deallocate a device based PCG solver.
type(ksp_monitor_t) function, dimension(3) cg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, blstx, blsty, blstz, gs_h, niter)
Standard PCG coupled solve.
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_add2s2(a_d, b_d, c1, n, strm)
Vector addition with scalar multiplication (multiplication on first argument)
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
real(kind=rp) function, public device_glsc3(a_d, b_d, c_d, n, strm)
Weighted inner product .
Device abstraction, common interface for various accelerators.
subroutine, public device_event_sync(event)
Synchronize an event.
subroutine, public device_event_destroy(event)
Destroy a device event.
subroutine, public device_event_create(event, flags)
Create a device event queue.
Implements the base abstract type for Krylov solvers plus helper types.
integer, parameter, public ksp_max_iter
Maximum number of iters.
integer, parameter, public rp
Global precision used in computations.
Base type for a matrix-vector product providing .
A list of allocatable `bc_t`. Follows the standard interface of lists.
Device based preconditioned conjugate gradient method.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Type for storing initial and final residuals in a Krylov solver.
Base abstract type for a canonical Krylov method, solving .
Defines a canonical Krylov preconditioner.