53 use mpi_f08,
only : mpi_iallreduce, mpi_status, &
54 mpi_sum, mpi_in_place, mpi_request, mpi_wait
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 :: p(:)
65 real(kind=
rp),
allocatable :: q(:)
66 real(kind=
rp),
allocatable :: r(:)
67 real(kind=
rp),
allocatable :: s(:)
68 real(kind=
rp),
allocatable :: u(:,:)
69 real(kind=
rp),
allocatable :: w(:)
70 real(kind=
rp),
allocatable :: z(:)
71 real(kind=
rp),
allocatable :: mi(:)
72 real(kind=
rp),
allocatable :: ni(:)
73 real(kind=
rp),
allocatable :: alpha(:)
74 real(kind=
rp),
allocatable :: beta(:)
75 type(c_ptr) :: p_d = c_null_ptr
76 type(c_ptr) :: q_d = c_null_ptr
77 type(c_ptr) :: r_d = c_null_ptr
78 type(c_ptr) :: s_d = c_null_ptr
79 type(c_ptr) :: u_d_d = c_null_ptr
80 type(c_ptr) :: w_d = c_null_ptr
81 type(c_ptr) :: z_d = c_null_ptr
82 type(c_ptr) :: mi_d = c_null_ptr
83 type(c_ptr) :: ni_d = c_null_ptr
84 type(c_ptr) :: alpha_d = c_null_ptr
85 type(c_ptr) :: beta_d = c_null_ptr
86 type(c_ptr),
allocatable :: u_d(:)
87 type(c_ptr) :: gs_event = c_null_ptr
98 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n) &
99 bind(c, name =
'cuda_pipecg_vecops')
100 use,
intrinsic :: iso_c_binding
103 type(c_ptr),
value :: p_d, q_d, r_d, s_d, u_d1, u_d2
104 type(c_ptr),
value :: w_d, ni_d, mi_d, z_d, mult_d
106 real(c_rp) :: alpha, beta, reduction(3)
113 bind(c, name =
'cuda_cg_update_xp')
114 use,
intrinsic :: iso_c_binding
116 type(c_ptr),
value :: x_d, p_d, u_d_d, alpha, beta
117 integer(c_int) :: p_cur, n, p_space
123 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n) &
124 bind(c, name =
'hip_pipecg_vecops')
125 use,
intrinsic :: iso_c_binding
128 type(c_ptr),
value :: p_d, q_d, r_d, s_d, u_d1, u_d2
129 type(c_ptr),
value :: w_d, ni_d, mi_d, z_d, mult_d
131 real(c_rp) :: alpha, beta, reduction(3)
138 bind(c, name =
'hip_cg_update_xp')
139 use,
intrinsic :: iso_c_binding
141 type(c_ptr),
value :: x_d, p_d, u_d_d, alpha, beta
142 integer(c_int) :: p_cur, n, p_space
150 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n)
151 type(c_ptr),
value :: p_d, q_d, r_d, s_d, u_d1, u_d2
152 type(c_ptr),
value :: w_d, ni_d, mi_d, z_d, mult_d
154 real(c_rp) :: alpha, beta, reduction(3)
157 s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, &
161 s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, &
164 call neko_error(
'No device backend configured')
170 use,
intrinsic :: iso_c_binding
171 type(c_ptr),
value :: x_d, p_d, u_d_d, alpha, beta
172 integer(c_int) :: p_cur, n, p_space
178 call neko_error(
'No device backend configured')
186 class(
pc_t),
optional,
intent(in),
target :: M
187 integer,
intent(in) :: n
188 integer,
intent(in) :: max_iter
189 real(kind=
rp),
optional,
intent(in) :: rel_tol
190 real(kind=
rp),
optional,
intent(in) :: abs_tol
191 logical,
optional,
intent(in) :: monitor
193 integer(c_size_t) :: u_size
226 this%u_d(i) = c_null_ptr
232 ptr = c_loc(this%u_d)
236 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
237 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
238 else if (
present(rel_tol) .and.
present(abs_tol))
then
239 call this%ksp_init(max_iter, rel_tol, abs_tol)
240 else if (
present(monitor) .and.
present(abs_tol))
then
241 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
242 else if (
present(rel_tol) .and.
present(monitor))
then
243 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
244 else if (
present(rel_tol))
then
245 call this%ksp_init(max_iter, rel_tol = rel_tol)
246 else if (
present(abs_tol))
then
247 call this%ksp_init(max_iter, abs_tol = abs_tol)
248 else if (
present(monitor))
then
249 call this%ksp_init(max_iter, monitor = monitor)
251 call this%ksp_init(max_iter)
265 if (
allocated(this%p))
then
266 if (c_associated(this%p_d))
then
271 if (
allocated(this%q))
then
272 if (c_associated(this%q_d))
then
277 if (
allocated(this%r))
then
278 if (c_associated(this%r_d))
then
283 if (
allocated(this%s))
then
284 if (c_associated(this%s_d))
then
289 if (
allocated(this%u))
then
290 if (
allocated(this%u_d))
then
292 if (c_associated(this%u_d(i)))
then
299 if (
allocated(this%u_d))
then
302 if (
allocated(this%w))
then
303 if (c_associated(this%w_d))
then
308 if (
allocated(this%z))
then
309 if (c_associated(this%z_d))
then
314 if (
allocated(this%mi))
then
315 if (c_associated(this%mi_d))
then
320 if (
allocated(this%ni))
then
321 if (c_associated(this%ni_d))
then
326 if (
allocated(this%alpha))
then
327 if (c_associated(this%alpha_d))
then
330 deallocate(this%alpha)
332 if (
allocated(this%beta))
then
333 if (c_associated(this%beta_d))
then
336 deallocate(this%beta)
339 if (c_associated(this%u_d_d))
then
345 if (c_associated(this%gs_event))
then
353 niter)
result(ksp_results)
355 class(
ax_t),
intent(in) :: ax
356 type(
field_t),
intent(inout) :: x
357 integer,
intent(in) :: n
358 real(kind=
rp),
dimension(n),
intent(in) :: f
359 type(
coef_t),
intent(inout) :: coef
361 type(
gs_t),
intent(inout) :: gs_h
363 integer,
optional,
intent(in) :: niter
364 integer :: iter, max_iter, ierr, p_cur, p_prev, u_prev
365 real(kind=
rp) :: rnorm, rtr, reduction(3), norm_fac
366 real(kind=
rp) :: gamma1, gamma2, delta
367 real(kind=
rp) :: tmp1, tmp2, tmp3
368 type(mpi_request) :: request
369 type(mpi_status) :: status
373 if (
present(niter))
then
376 max_iter = this%max_iter
378 norm_fac = 1.0_rp / sqrt(coef%volume)
380 associate(p => this%p, q => this%q, r => this%r, s => this%s, &
381 u => this%u, w => this%w, z => this%z, mi => this%mi, ni => this%ni, &
382 alpha => this%alpha, beta => this%beta, &
383 alpha_d => this%alpha_d, beta_d => this%beta_d, &
384 p_d => this%p_d, q_d => this%q_d, r_d => this%r_d, &
385 s_d => this%s_d, u_d => this%u_d, u_d_d => this%u_d_d, &
386 w_d => this%w_d, z_d => this%z_d, mi_d => this%mi_d, ni_d => this%ni_d)
399 call this%M%solve(u(1, u_prev), r, n)
400 call ax%compute(w, u(1, u_prev), coef, x%msh, x%Xh)
401 call gs_h%op(w, n, gs_op_add, this%gs_event)
403 call bc_projector%apply(w, n)
406 rnorm = sqrt(rtr)*norm_fac
407 ksp_results%res_start = rnorm
408 ksp_results%res_final = rnorm
410 if (
abscmp(rnorm, 0.0_rp))
then
411 ksp_results%converged = .true.
426 call this%monitor_start(
'PipeCG')
427 do iter = 1, max_iter
428 call mpi_iallreduce(mpi_in_place, reduction, 3, &
431 call this%M%solve(mi, w, n)
432 call ax%compute(ni, mi, coef, x%msh, x%Xh)
433 call gs_h%op(ni, n, gs_op_add, this%gs_event)
435 call bc_projector%apply(ni, n)
437 call mpi_wait(request, status, ierr)
439 gamma1 = reduction(1)
443 rnorm = sqrt(rtr)*norm_fac
444 call this%monitor_iter(iter, rnorm)
445 if (rnorm .lt. this%abs_tol)
exit
448 if (iter .gt. 1)
then
449 beta(p_cur) = gamma1 / gamma2
451 gamma1 / (delta - (beta(p_cur) * gamma1/alpha(p_prev)))
454 alpha(p_cur) = gamma1/delta
458 s_d, u_d(u_prev), u_d(p_cur),&
460 mi_d, alpha(p_cur), beta(p_cur),&
461 coef%mult_d, reduction, n)
471 alpha(1) = alpha(p_cur)
472 beta(1) = beta(p_cur)
481 if ( p_cur .ne. 1)
then
488 call this%monitor_stop()
489 ksp_results%res_final = rnorm
490 ksp_results%iter = iter
491 ksp_results%converged = this%is_converged(iter, rnorm)
499 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
501 class(ax_t),
intent(in) :: ax
502 type(field_t),
intent(inout) :: x
503 type(field_t),
intent(inout) :: y
504 type(field_t),
intent(inout) :: z
505 integer,
intent(in) :: n
506 real(kind=rp),
dimension(n),
intent(in) :: fx
507 real(kind=rp),
dimension(n),
intent(in) :: fy
508 real(kind=rp),
dimension(n),
intent(in) :: fz
509 type(coef_t),
intent(inout) :: coef
510 class(vector_bc_projector_t),
intent(inout) :: bc_projector
511 type(gs_t),
intent(inout) :: gs_h
512 type(ksp_monitor_t),
dimension(3) :: ksp_results
513 integer,
optional,
intent(in) :: niter
514 type(scalar_bc_projector_t),
pointer :: bc_x, bc_y, bc_z
516 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
517 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
518 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
519 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)
Copy data between host and device (or device and device)
Unmap a Fortran array from a device (deassociate and free)
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_vlsc3(u_d, v_d, w_d, n, strm)
Compute multiplication sum .
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.
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.
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.
Defines a pipelined Conjugate Gradient methods.
subroutine device_pipecg_vecops(p_d, q_d, r_d, s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction, n)
subroutine pipecg_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a pipelined PCG solver.
type(ksp_monitor_t) function pipecg_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Pipelined PCG solve.
subroutine device_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, p_cur, p_space, n)
integer, parameter device_pipecg_p_space
type(ksp_monitor_t) function, dimension(3) pipecg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Pipelined PCG coupled solve.
subroutine pipecg_device_free(this)
Deallocate a pipelined PCG solver.
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.
void hip_cg_update_xp(void *x, void *p, void *u, void *alpha, void *beta, int *p_cur, int *p_space, int *n)
void hip_pipecg_vecops(void *p, void *q, void *r, void *s, void *u1, void *u2, void *w, void *z, void *ni, void *mi, real *alpha, real *beta, void *mult, real *reduction, int *n)
Base type for a matrix-vector product providing .
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 .
Pipelined preconditioned conjugate gradient method.
Defines a canonical Krylov preconditioner.
Projector for scalar boundary conditions.
Abstract type for resolving vector boundary conditions.