52 use mpi_f08,
only : mpi_in_place, mpi_allreduce, &
55 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, &
56 c_associated, c_size_t, c_sizeof, c_int, c_loc
64 real(kind=
rp),
allocatable :: w1(:)
65 real(kind=
rp),
allocatable :: w2(:)
66 real(kind=
rp),
allocatable :: w3(:)
67 real(kind=
rp),
allocatable :: r1(:)
68 real(kind=
rp),
allocatable :: r2(:)
69 real(kind=
rp),
allocatable :: r3(:)
70 real(kind=
rp),
allocatable :: z1(:)
71 real(kind=
rp),
allocatable :: z2(:)
72 real(kind=
rp),
allocatable :: z3(:)
73 real(kind=
rp),
allocatable :: tmp(:)
74 real(kind=
rp),
allocatable :: p1(:,:)
75 real(kind=
rp),
allocatable :: p2(:,:)
76 real(kind=
rp),
allocatable :: p3(:,:)
77 real(kind=
rp),
allocatable :: alpha(:)
78 type(c_ptr) :: w1_d = c_null_ptr
79 type(c_ptr) :: w2_d = c_null_ptr
80 type(c_ptr) :: w3_d = c_null_ptr
81 type(c_ptr) :: r1_d = c_null_ptr
82 type(c_ptr) :: r2_d = c_null_ptr
83 type(c_ptr) :: r3_d = c_null_ptr
84 type(c_ptr) :: z1_d = c_null_ptr
85 type(c_ptr) :: z2_d = c_null_ptr
86 type(c_ptr) :: z3_d = c_null_ptr
87 type(c_ptr) :: alpha_d = c_null_ptr
88 type(c_ptr) :: p1_d_d = c_null_ptr
89 type(c_ptr) :: p2_d_d = c_null_ptr
90 type(c_ptr) :: p3_d_d = c_null_ptr
91 type(c_ptr) :: tmp_d = c_null_ptr
92 type(c_ptr),
allocatable :: p1_d(:)
93 type(c_ptr),
allocatable :: p2_d(:)
94 type(c_ptr),
allocatable :: p3_d(:)
95 type(c_ptr) :: gs_event1 = c_null_ptr
96 type(c_ptr) :: gs_event2 = c_null_ptr
97 type(c_ptr) :: gs_event3 = c_null_ptr
108 b1_d, b2_d, b3_d, tmp_d, n) bind(c, name = 'cuda_fusedcg_cpld_part1')
109 use,
intrinsic :: iso_c_binding
112 type(c_ptr),
value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d
119 po1_d, po2_d, po3_d, beta, n) &
120 bind(c, name =
'cuda_fusedcg_cpld_update_p')
121 use,
intrinsic :: iso_c_binding
124 type(c_ptr),
value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
125 type(c_ptr),
value :: po1_d, po2_d, po3_d
133 alpha, p_cur, n) bind(c, name = 'cuda_fusedcg_cpld_update_x')
134 use,
intrinsic :: iso_c_binding
136 type(c_ptr),
value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
137 integer(c_int) :: p_cur, n
143 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n) &
144 bind(c, name =
'cuda_fusedcg_cpld_part2')
145 use,
intrinsic :: iso_c_binding
148 type(c_ptr),
value :: a1_d, a2_d, a3_d, b_d
149 type(c_ptr),
value :: c1_d, c2_d, c3_d, alpha_d
151 integer(c_int) :: n, p_cur
157 b1_d, b2_d, b3_d, tmp_d, n) &
158 bind(c, name =
'hip_fusedcg_cpld_part1')
159 use,
intrinsic :: iso_c_binding
162 type(c_ptr),
value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d
169 po1_d, po2_d, po3_d, beta, n) &
170 bind(c, name =
'hip_fusedcg_cpld_update_p')
171 use,
intrinsic :: iso_c_binding
174 type(c_ptr),
value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
175 type(c_ptr),
value :: po1_d, po2_d, po3_d
183 alpha, p_cur, n) bind(c, name = 'hip_fusedcg_cpld_update_x')
184 use,
intrinsic :: iso_c_binding
186 type(c_ptr),
value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
187 integer(c_int) :: p_cur, n
193 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n) &
194 bind(c, name =
'hip_fusedcg_cpld_part2')
195 use,
intrinsic :: iso_c_binding
198 type(c_ptr),
value :: a1_d, a2_d, a3_d, b_d
199 type(c_ptr),
value :: c1_d, c2_d, c3_d, alpha_d
201 integer(c_int) :: n, p_cur
209 b1_d, b2_d, b3_d, tmp_d, n)
210 type(c_ptr),
value :: a1_d, a2_d, a3_d, b1_d, b2_d, b3_d
211 type(c_ptr),
value :: tmp_d
218 call neko_error(
'No device backend configured')
223 po1_d, po2_d, po3_d, beta, n)
224 type(c_ptr),
value :: p1_d, p2_d, p3_d, z1_d, z2_d, z3_d
225 type(c_ptr),
value :: po1_d, po2_d, po3_d
230 po1_d, po2_d, po3_d, beta, n)
233 po1_d, po2_d, po3_d, beta, n)
235 call neko_error(
'No device backend configured')
240 p1_d, p2_d, p3_d, alpha, p_cur, n)
241 type(c_ptr),
value :: x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha
242 integer(c_int) :: p_cur, n
245 p1_d, p2_d, p3_d, alpha, p_cur, n)
248 p1_d, p2_d, p3_d, alpha, p_cur, n)
250 call neko_error(
'No device backend configured')
255 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
result(res)
256 type(c_ptr),
value :: a1_d, a2_d, a3_d, b_d
257 type(c_ptr),
value :: c1_d, c2_d, c3_d, alpha_d
264 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
267 c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
269 call neko_error(
'No device backend configured')
272#ifndef HAVE_DEVICE_MPI
274 call mpi_allreduce(mpi_in_place, res, 1, &
283 rel_tol, abs_tol, monitor)
285 class(
pc_t),
optional,
intent(in),
target :: M
286 integer,
intent(in) :: n
287 integer,
intent(in) :: max_iter
288 real(kind=
rp),
optional,
intent(in) :: rel_tol
289 real(kind=
rp),
optional,
intent(in) :: abs_tol
290 logical,
optional,
intent(in) :: monitor
292 integer(c_size_t) :: p_size
306 allocate(this%tmp(n))
331 this%p1_d(i) = c_null_ptr
332 call device_map(this%p1(:,i), this%p1_d(i), n)
334 this%p2_d(i) = c_null_ptr
335 call device_map(this%p2(:,i), this%p2_d(i), n)
337 this%p3_d(i) = c_null_ptr
338 call device_map(this%p3(:,i), this%p3_d(i), n)
345 ptr = c_loc(this%p1_d)
348 ptr = c_loc(this%p2_d)
351 ptr = c_loc(this%p3_d)
354 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
355 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
356 else if (
present(rel_tol) .and.
present(abs_tol))
then
357 call this%ksp_init(max_iter, rel_tol, abs_tol)
358 else if (
present(monitor) .and.
present(abs_tol))
then
359 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
360 else if (
present(rel_tol) .and.
present(monitor))
then
361 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
362 else if (
present(rel_tol))
then
363 call this%ksp_init(max_iter, rel_tol = rel_tol)
364 else if (
present(abs_tol))
then
365 call this%ksp_init(max_iter, abs_tol = abs_tol)
366 else if (
present(monitor))
then
367 call this%ksp_init(max_iter, monitor = monitor)
369 call this%ksp_init(max_iter)
385 if (
allocated(this%w1))
then
386 if (c_associated(this%w1_d))
then
392 if (
allocated(this%w2))
then
393 if (c_associated(this%w2_d))
then
399 if (
allocated(this%w3))
then
400 if (c_associated(this%w3_d))
then
406 if (
allocated(this%r1))
then
407 if (c_associated(this%r1_d))
then
413 if (
allocated(this%r2))
then
414 if (c_associated(this%r2_d))
then
420 if (
allocated(this%r3))
then
421 if (c_associated(this%r3_d))
then
427 if (
allocated(this%z1))
then
428 if (c_associated(this%z1_d))
then
434 if (
allocated(this%z2))
then
435 if (c_associated(this%z2_d))
then
441 if (
allocated(this%z3))
then
442 if (c_associated(this%z3_d))
then
448 if (
allocated(this%tmp))
then
449 if (c_associated(this%tmp_d))
then
455 if (
allocated(this%alpha))
then
456 if (c_associated(this%alpha_d))
then
459 deallocate(this%alpha)
462 if (
allocated(this%p1))
then
463 if (
allocated(this%p1_d))
then
465 if (c_associated(this%p1_d(i)))
then
473 if (
allocated(this%p2))
then
474 if (
allocated(this%p2_d))
then
476 if (c_associated(this%p2_d(i)))
then
484 if (
allocated(this%p3))
then
485 if (
allocated(this%p3_d))
then
487 if (c_associated(this%p3_d(i)))
then
495 if (
allocated(this%p1_d))
then
496 deallocate(this%p1_d)
499 if (
allocated(this%p2_d))
then
500 deallocate(this%p2_d)
503 if (
allocated(this%p3_d))
then
504 deallocate(this%p3_d)
507 if (c_associated(this%p1_d_d))
then
511 if (c_associated(this%p2_d_d))
then
515 if (c_associated(this%p3_d_d))
then
521 if (c_associated(this%gs_event1))
then
525 if (c_associated(this%gs_event2))
then
529 if (c_associated(this%gs_event3))
then
537 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
539 class(
ax_t),
intent(in) :: ax
540 type(
field_t),
intent(inout) :: x
541 type(
field_t),
intent(inout) :: y
542 type(
field_t),
intent(inout) :: z
543 integer,
intent(in) :: n
544 real(kind=
rp),
dimension(n),
intent(in) :: fx
545 real(kind=
rp),
dimension(n),
intent(in) :: fy
546 real(kind=
rp),
dimension(n),
intent(in) :: fz
547 type(
coef_t),
intent(inout) :: coef
549 type(
gs_t),
intent(inout) :: gs_h
551 integer,
optional,
intent(in) :: niter
552 integer :: iter, max_iter, ierr, i, p_cur, p_prev
553 real(kind=
rp) :: rnorm, rtr, norm_fac, rtz1, rtz2
554 real(kind=
rp) :: pap, beta
563 if (
present(niter))
then
568 norm_fac = 1.0_rp / sqrt(coef%volume)
570 associate(w1 => this%w1, w2 => this%w2, w3 => this%w3, r1 => this%r1, &
571 r2 => this%r2, r3 => this%r3, p1 => this%p1, p2 => this%p2, &
572 p3 => this%p3, z1 => this%z1, z2 => this%z2, z3 => this%z3, &
573 tmp_d => this%tmp_d, alpha => this%alpha, alpha_d => this%alpha_d, &
574 w1_d => this%w1_d, w2_d => this%w2_d, w3_d => this%w3_d, &
575 r1_d => this%r1_d, r2_d => this%r2_d, r3_d => this%r3_d, &
576 z1_d => this%z1_d, z2_d => this%z2_d, z3_d => this%z3_d, &
577 p1_d => this%p1_d, p2_d => this%p2_d, p3_d => this%p3_d, &
578 p1_d_d => this%p1_d_d, p2_d_d => this%p2_d_d, p3_d_d => this%p3_d_d)
596 r2_d, r3_d, tmp_d, n)
600 rnorm = sqrt(rtr)*norm_fac
601 ksp_results%res_start = rnorm
602 ksp_results%res_final = rnorm
603 ksp_results(1)%iter = 0
604 ksp_results(2:3)%iter = -1
605 if (
abscmp(rnorm, 0.0_rp))
then
606 ksp_results%converged = .true.
609 call this%monitor_start(
'fcpldCG')
610 do iter = 1, max_iter
611 call this%M%solve(z1, r1, n)
612 call this%M%solve(z2, r2, n)
613 call this%M%solve(z3, r3, n)
616 r1_d, r2_d, r3_d, tmp_d, n)
620 if (iter .eq. 1) beta = 0.0_rp
623 z1_d, z2_d, z3_d, p1_d(p_prev), p2_d(p_prev), p3_d(p_prev), beta, n)
625 call ax%compute_vector(w1, w2, w3, &
626 p1(1, p_cur), p2(1, p_cur), p3(1, p_cur), coef, x%msh, x%Xh)
629 call gs_h%op(w1, w2, w3, n, gs_op_add, this%gs_event1)
631 call bc_projector%apply(w1, w2, w3, n)
635 p2_d(p_cur), p3_d(p_cur), tmp_d, n)
639 alpha(p_cur) = rtz1 / pap
641 w1_d, w2_d, w3_d, alpha_d, alpha(p_cur), p_cur, n)
642 rnorm = sqrt(rtr)*norm_fac
643 call this%monitor_iter(iter, rnorm)
645 (rnorm .lt. this%abs_tol) .or. iter .eq. max_iter)
then
647 p1_d_d, p2_d_d, p3_d_d, alpha_d, p_cur, n)
650 if (rnorm .lt. this%abs_tol)
exit
656 call this%monitor_stop()
657 ksp_results%res_final = rnorm
658 ksp_results%iter = iter
659 ksp_results%converged = this%is_converged(iter, rnorm)
667 gs_h, niter)
result(ksp_results)
669 class(
ax_t),
intent(in) :: ax
670 type(
field_t),
intent(inout) :: x
671 integer,
intent(in) :: n
672 real(kind=
rp),
dimension(n),
intent(in) :: f
673 type(
coef_t),
intent(inout) :: coef
675 type(
gs_t),
intent(inout) :: gs_h
677 integer,
optional,
intent(in) :: niter
680 call neko_error(
'The cpldcg solver is only defined for coupled solves')
682 ksp_results%res_final = 0.0
684 ksp_results%converged = .false.
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
void hip_fusedcg_cpld_update_x(void *x1, void *x2, void *x3, void *p1, void *p2, void *p3, void *alpha, int *p_cur, int *n)
void hip_fusedcg_cpld_update_p(void *p1, void *p2, void *p3, void *z1, void *z2, void *z3, void *po1, void *po2, void *po3, real *beta, int *n)
real hip_fusedcg_cpld_part2(void *a1, void *a2, void *a3, void *b, void *c1, void *c2, void *c3, void *alpha_d, real *alpha, int *p_cur, int *n)
void hip_fusedcg_cpld_part1(void *a1, void *a2, void *a3, void *b1, void *b2, void *b3, void *tmp, int *n)
Return the device pointer for an associated Fortran array.
Map a Fortran array to a device (allocate and associate)
Copy data between host and device (or device and device)
Unmap a Fortran array from a device (deassociate and free)
Apply cyclic boundary condition to a vector field.
Defines a Matrix-vector product.
type(mpi_datatype), public mpi_real_precision
MPI type for working precision of REAL types.
integer, public pe_size
MPI size of communicator.
type(mpi_comm), public neko_comm
MPI communicator.
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_glsc2(a_d, b_d, n, strm)
Weighted inner product .
Device abstraction, common interface for various accelerators.
subroutine, public device_event_sync(event)
Synchronize an event.
integer, parameter, public host_to_device
subroutine, public device_free(x_d)
Deallocate memory on the device.
subroutine, public device_event_destroy(event)
Destroy a device event.
subroutine, public device_alloc(x_d, s)
Allocate memory on the device.
subroutine, public device_event_create(event, flags)
Create a device event queue.
Defines a fused Conjugate Gradient method for accelerators.
subroutine device_fusedcg_cpld_update_x(x1_d, x2_d, x3_d, p1_d, p2_d, p3_d, alpha, p_cur, n)
type(ksp_monitor_t) function, dimension(3) fusedcg_cpld_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Pipelined PCG solve coupled solve.
type(ksp_monitor_t) function fusedcg_cpld_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Pipelined PCG solve.
subroutine fusedcg_cpld_device_free(this)
Deallocate a pipelined PCG solver.
subroutine device_fusedcg_cpld_update_p(p1_d, p2_d, p3_d, z1_d, z2_d, z3_d, po1_d, po2_d, po3_d, beta, n)
subroutine fusedcg_cpld_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a fused PCG solver.
real(kind=rp) function device_fusedcg_cpld_part2(a1_d, a2_d, a3_d, b_d, c1_d, c2_d, c3_d, alpha_d, alpha, p_cur, n)
integer, parameter device_fusedcg_cpld_p_space
subroutine device_fusedcg_cpld_part1(a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, tmp_d, n)
Implements the base abstract type for Krylov solvers plus helper types.
integer, parameter, public ksp_max_iter
Maximum number of iters.
real(kind=rp) function, public glsc3(a, b, c, n)
Weighted inner product .
subroutine, public copy(a, b, n)
Copy a vector .
subroutine, public rzero(a, n)
Zero a real vector.
integer, parameter, public c_rp
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 .
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Fused preconditioned conjugate gradient method.
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.