50 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
56 real(kind=
rp),
allocatable :: w(:)
57 real(kind=
rp),
allocatable :: r(:)
58 real(kind=
rp),
allocatable :: p(:)
59 real(kind=
rp),
allocatable :: z(:)
60 type(c_ptr) :: w_d = c_null_ptr
61 type(c_ptr) :: r_d = c_null_ptr
62 type(c_ptr) :: p_d = c_null_ptr
63 type(c_ptr) :: z_d = c_null_ptr
64 type(c_ptr) :: gs_event = c_null_ptr
77 class(
pc_t),
optional,
intent(in),
target :: M
78 integer,
intent(in) :: n
79 integer,
intent(in) :: max_iter
80 real(kind=
rp),
optional,
intent(in) :: rel_tol
81 real(kind=
rp),
optional,
intent(in) :: abs_tol
82 logical,
optional,
intent(in) :: monitor
100 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
101 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
102 else if (
present(rel_tol) .and.
present(abs_tol))
then
103 call this%ksp_init(max_iter, rel_tol, abs_tol)
104 else if (
present(monitor) .and.
present(abs_tol))
then
105 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
106 else if (
present(rel_tol) .and.
present(monitor))
then
107 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
108 else if (
present(rel_tol))
then
109 call this%ksp_init(max_iter, rel_tol = rel_tol)
110 else if (
present(abs_tol))
then
111 call this%ksp_init(max_iter, abs_tol = abs_tol)
112 else if (
present(monitor))
then
113 call this%ksp_init(max_iter, monitor = monitor)
115 call this%ksp_init(max_iter)
127 if (
allocated(this%w))
then
128 if (c_associated(this%w_d))
then
134 if (
allocated(this%r))
then
135 if (c_associated(this%r_d))
then
141 if (
allocated(this%p))
then
142 if (c_associated(this%p_d))
then
148 if (
allocated(this%z))
then
149 if (c_associated(this%z_d))
then
157 if (c_associated(this%gs_event))
then
167 class(
ax_t),
intent(in) :: ax
168 type(
field_t),
intent(inout) :: x
169 integer,
intent(in) :: n
170 real(kind=
rp),
dimension(n),
intent(in) :: f
171 type(
coef_t),
intent(inout) :: coef
173 type(
gs_t),
intent(inout) :: gs_h
175 integer,
optional,
intent(in) :: niter
176 real(kind=
rp),
parameter :: one = 1.0
177 real(kind=
rp),
parameter :: zero = 0.0
178 integer :: iter, max_iter
179 real(kind=
rp) :: rnorm, rtr, rtr0, rtz2, rtz1
180 real(kind=
rp) :: beta, pap, alpha, alphm, norm_fac
185 if (
present(niter))
then
188 max_iter = this%max_iter
190 norm_fac = one/sqrt(coef%volume)
198 rnorm = sqrt(rtr)*norm_fac
199 ksp_results%res_start = rnorm
200 ksp_results%res_final = rnorm
202 if (
abscmp(rnorm, zero))
then
203 ksp_results%converged = .true.
206 call this%monitor_start(
'CG')
207 do iter = 1, max_iter
208 call this%M%solve(this%z, this%r, n)
212 if (iter .eq. 1) beta = zero
215 call ax%compute(this%w, this%p, coef, x%msh, x%Xh)
216 call gs_h%op(this%w, n, gs_op_add, this%gs_event)
218 call bc_projector%apply(this%w, n)
228 if (iter .eq. 1) rtr0 = rtr
229 rnorm = sqrt(rtr)*norm_fac
230 call this%monitor_iter(iter, rnorm)
231 if (rnorm .lt. this%abs_tol)
then
235 call this%monitor_stop()
236 ksp_results%res_final = rnorm
237 ksp_results%iter = iter
238 ksp_results%converged = this%is_converged(iter, rnorm)
244 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
246 class(
ax_t),
intent(in) :: ax
247 type(
field_t),
intent(inout) :: x
248 type(
field_t),
intent(inout) :: y
249 type(
field_t),
intent(inout) :: z
250 integer,
intent(in) :: n
251 real(kind=
rp),
dimension(n),
intent(in) :: fx
252 real(kind=
rp),
dimension(n),
intent(in) :: fy
253 real(kind=
rp),
dimension(n),
intent(in) :: fz
254 type(
coef_t),
intent(inout) :: coef
256 type(
gs_t),
intent(inout) :: gs_h
258 integer,
optional,
intent(in) :: niter
262 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
263 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
264 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, 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, dimension(3) cg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard PCG coupled 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 cg_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Standard PCG 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.
Implements scalar_projector_t.
Implements boundary condition projectors for vector fields. Two types concrete types are provided: se...
subroutine, public vector_bc_projector_components(this, x, y, z)
Access the component scalar projectors from a segregated vector projector.
Base type for a matrix-vector product providing .
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.
Projector for scalar boundary conditions.
Abstract type for resolving vector boundary conditions.