53 use mpi_f08,
only : mpi_allreduce, mpi_in_place, mpi_sum
54 use,
intrinsic :: ieee_arithmetic, only : ieee_is_finite
55 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated, &
73 real(kind=
rp),
allocatable :: p(:)
75 real(kind=
rp),
allocatable :: p_hat(:)
77 real(kind=
rp),
allocatable :: r(:)
79 real(kind=
rp),
allocatable :: s(:)
81 real(kind=
rp),
allocatable :: s_hat(:)
83 real(kind=
rp),
allocatable :: t(:)
85 real(kind=
rp),
allocatable :: v(:)
86 type(c_ptr) :: p_d = c_null_ptr
87 type(c_ptr) :: p_hat_d = c_null_ptr
88 type(c_ptr) :: r_d = c_null_ptr
89 type(c_ptr) :: s_d = c_null_ptr
90 type(c_ptr) :: s_hat_d = c_null_ptr
91 type(c_ptr) :: t_d = c_null_ptr
92 type(c_ptr) :: v_d = c_null_ptr
93 type(c_ptr) :: gs_event = c_null_ptr
108 bind(c, name =
'hip_bicgstab_update_p')
109 use,
intrinsic :: iso_c_binding
112 type(c_ptr),
value :: p_d, r_d, v_d
113 real(
c_rp) :: beta, omega
120 bind(c, name =
'hip_bicgstab_product_and_norm')
121 use,
intrinsic :: iso_c_binding
124 type(c_ptr),
value :: a_d, b_d, mult_d
132 alpha, n) bind(c, name =
'hip_bicgstab_part1')
133 use,
intrinsic :: iso_c_binding
136 type(c_ptr),
value :: s_d, r_d, v_d, mult_d
144 t_d, f_d, mult_d, alpha, omega, res, n) &
145 bind(c, name =
'hip_bicgstab_part2')
146 use,
intrinsic :: iso_c_binding
149 type(c_ptr),
value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
150 type(c_ptr),
value :: mult_d
151 real(c_rp) :: alpha, omega
159 bind(c, name =
'cuda_bicgstab_update_p')
160 use,
intrinsic :: iso_c_binding
163 type(c_ptr),
value :: p_d, r_d, v_d
164 real(c_rp) :: beta, omega
171 bind(c, name =
'cuda_bicgstab_product_and_norm')
172 use,
intrinsic :: iso_c_binding
175 type(c_ptr),
value :: a_d, b_d, mult_d
182 real(c_xp) function cuda_bicgstab_part1(s_d, r_d, v_d, mult_d, &
183 alpha, n) bind(c, name =
'cuda_bicgstab_part1')
184 use,
intrinsic :: iso_c_binding
187 type(c_ptr),
value :: s_d, r_d, v_d, mult_d
195 t_d, f_d, mult_d, alpha, omega, res, n) &
196 bind(c, name =
'cuda_bicgstab_part2')
197 use,
intrinsic :: iso_c_binding
200 type(c_ptr),
value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
201 type(c_ptr),
value :: mult_d
202 real(c_rp) :: alpha, omega
210 strm) bind(c, name = 'opencl_bicgstab_update_p')
211 use,
intrinsic :: iso_c_binding
214 type(c_ptr),
value :: p_d, r_d, v_d, strm
215 real(c_rp) :: beta, omega
222 n, strm) bind(c, name = 'opencl_bicgstab_product_and_norm')
223 use,
intrinsic :: iso_c_binding
226 type(c_ptr),
value :: a_d, b_d, mult_d, strm
233 real(c_xp) function opencl_bicgstab_part1(s_d, r_d, v_d, mult_d, &
234 alpha, n, strm) bind(c, name =
'opencl_bicgstab_part1')
235 use,
intrinsic :: iso_c_binding
238 type(c_ptr),
value :: s_d, r_d, v_d, mult_d, strm
246 t_d, f_d, mult_d, alpha, omega, res, n, strm) &
247 bind(c, name =
'opencl_bicgstab_part2')
248 use,
intrinsic :: iso_c_binding
251 type(c_ptr),
value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
252 type(c_ptr),
value :: mult_d, strm
253 real(c_rp) :: alpha, omega
260 subroutine metal_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
261 strm) bind(c, name = 'metal_bicgstab_update_p')
262 use,
intrinsic :: iso_c_binding
265 type(c_ptr),
value :: p_d, r_d, v_d, strm
266 real(c_rp) :: beta, omega
268 end subroutine metal_bicgstab_update_p
272 subroutine metal_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n, &
273 strm) bind(c, name = 'metal_bicgstab_product_and_norm')
274 use,
intrinsic :: iso_c_binding
277 type(c_ptr),
value :: a_d, b_d, mult_d, strm
280 end subroutine metal_bicgstab_product_and_norm
284 real(c_rp) function metal_bicgstab_part1(s_d, r_d, v_d, mult_d, &
285 alpha, n, strm) bind(c, name =
'metal_bicgstab_part1')
286 use,
intrinsic :: iso_c_binding
289 type(c_ptr),
value :: s_d, r_d, v_d, mult_d, strm
292 end function metal_bicgstab_part1
296 subroutine metal_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, &
297 t_d, f_d, mult_d, alpha, omega, res, n, strm) &
298 bind(c, name =
'metal_bicgstab_part2')
299 use,
intrinsic :: iso_c_binding
302 type(c_ptr),
value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
303 type(c_ptr),
value :: mult_d, strm
304 real(c_rp) :: alpha, omega, res(2)
306 end subroutine metal_bicgstab_part2
320 type(c_ptr) :: p_d, r_d, v_d
321 real(kind=rp) :: beta, omega
332 call metal_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
335 call neko_error(
'No device backend configured')
353 real(kind=rp),
intent(out) :: product
354 real(kind=rp),
intent(out) :: norm_squared
355 type(c_ptr) :: a_d, b_d, mult_d
357 real(kind=xp) :: res_xp(2)
360 real(kind=rp) :: res(2)
372 call metal_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n, &
374 res_xp =
real(res, kind=xp)
376 call neko_error(
'No device backend configured')
379#ifndef HAVE_DEVICE_MPI
380 if (pe_size .gt. 1)
then
381 call mpi_allreduce(mpi_in_place, res_xp, 2, &
382 mpi_extra_precision, mpi_sum, neko_comm, ierr)
385 product =
real(res_xp(1), kind=rp)
386 norm_squared =
real(res_xp(2), kind=rp)
400 type(c_ptr) :: s_d, r_d, v_d, mult_d
401 real(kind=rp) :: alpha
404 real(kind=xp) :: res_xp
416 res = metal_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n, &
418 res_xp =
real(res, kind=xp)
420 call neko_error(
'No device backend configured')
423#ifndef HAVE_DEVICE_MPI
424 if (pe_size .gt. 1)
then
425 call mpi_allreduce(mpi_in_place, res_xp, 1, &
426 mpi_extra_precision, mpi_sum, neko_comm, ierr)
429 res =
real(res_xp, kind=rp)
453 s_d, t_d, f_d, mult_d, alpha, omega, n)
454 real(kind=rp),
intent(out) :: rtr
455 real(kind=rp),
intent(out) :: rho
456 type(c_ptr) :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, mult_d
457 real(kind=rp) :: alpha, omega
459 real(kind=xp) :: res_xp(2)
462 real(kind=rp) :: res(2)
468 mult_d, alpha, omega, res_xp, n)
471 mult_d, alpha, omega, res_xp, n)
474 mult_d, alpha, omega, res_xp, n, glb_cmd_queue)
476 call metal_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, &
477 mult_d, alpha, omega, res, n, glb_cmd_queue)
478 res_xp =
real(res, kind=xp)
480 call neko_error(
'No device backend configured')
483#ifndef HAVE_DEVICE_MPI
484 if (pe_size .gt. 1)
then
485 call mpi_allreduce(mpi_in_place, res_xp, 2, &
486 mpi_extra_precision, mpi_sum, neko_comm, ierr)
489 rtr =
real(res_xp(1), kind=rp)
490 rho =
real(res_xp(2), kind=rp)
505 class(pc_t),
optional,
intent(in),
target :: M
506 integer,
intent(in) :: n
507 integer,
intent(in) :: max_iter
508 real(kind=rp),
optional,
intent(in) :: rel_tol
509 real(kind=rp),
optional,
intent(in) :: abs_tol
510 logical,
optional,
intent(in) :: monitor
515 allocate(this%p_hat(n))
518 allocate(this%s_hat(n))
522 call device_map(this%p, this%p_d, n)
523 call device_map(this%p_hat, this%p_hat_d, n)
524 call device_map(this%r, this%r_d, n)
525 call device_map(this%s, this%s_d, n)
526 call device_map(this%s_hat, this%s_hat_d, n)
527 call device_map(this%t, this%t_d, n)
528 call device_map(this%v, this%v_d, n)
534 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
535 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
536 else if (
present(rel_tol) .and.
present(abs_tol))
then
537 call this%ksp_init(max_iter, rel_tol, abs_tol)
538 else if (
present(monitor) .and.
present(abs_tol))
then
539 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
540 else if (
present(rel_tol) .and.
present(monitor))
then
541 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
542 else if (
present(rel_tol))
then
543 call this%ksp_init(max_iter, rel_tol = rel_tol)
544 else if (
present(abs_tol))
then
545 call this%ksp_init(max_iter, abs_tol = abs_tol)
546 else if (
present(monitor))
then
547 call this%ksp_init(max_iter, monitor = monitor)
549 call this%ksp_init(max_iter)
552 call device_event_create(this%gs_event, 2)
562 if (
allocated(this%v))
then
563 if (c_associated(this%v_d))
then
564 call device_unmap(this%v, this%v_d)
569 if (
allocated(this%r))
then
570 if (c_associated(this%r_d))
then
571 call device_unmap(this%r, this%r_d)
576 if (
allocated(this%t))
then
577 if (c_associated(this%t_d))
then
578 call device_unmap(this%t, this%t_d)
583 if (
allocated(this%p))
then
584 if (c_associated(this%p_d))
then
585 call device_unmap(this%p, this%p_d)
590 if (
allocated(this%p_hat))
then
591 if (c_associated(this%p_hat_d))
then
592 call device_unmap(this%p_hat, this%p_hat_d)
594 deallocate(this%p_hat)
597 if (
allocated(this%s))
then
598 if (c_associated(this%s_d))
then
599 call device_unmap(this%s, this%s_d)
604 if (
allocated(this%s_hat))
then
605 if (c_associated(this%s_hat_d))
then
606 call device_unmap(this%s_hat, this%s_hat_d)
608 deallocate(this%s_hat)
613 if (c_associated(this%gs_event))
then
614 call device_event_destroy(this%gs_event)
633 gs_h, niter)
result(ksp_results)
635 class(ax_t),
intent(in) :: ax
636 type(field_t),
intent(inout) :: x
637 integer,
intent(in) :: n
638 real(kind=rp),
dimension(n),
intent(in) :: f
639 type(coef_t),
intent(inout) :: coef
640 class(scalar_bc_projector_t),
intent(inout) :: bc_projector
641 type(gs_t),
intent(inout) :: gs_h
642 type(ksp_monitor_t) :: ksp_results
643 integer,
optional,
intent(in) :: niter
644 integer :: iter, max_iter
645 real(kind=rp) :: rnorm, rtr, norm_fac, gamma
646 real(kind=rp) :: r_norm, s_norm, shadow_norm, t_norm, v_norm
648 real(kind=rp) :: sts, ftv, vtv, stt, ttt
649 real(kind=rp) :: beta, alpha, omega, rho_1, rho_2, rho_next
652 f_d = device_get_ptr(f)
654 if (
present(niter))
then
657 max_iter = this%max_iter
659 norm_fac = 1.0_rp / sqrt(coef%volume)
661 associate(r_d => this%r_d, t_d => this%t_d, s_d => this%s_d, &
662 v_d => this%v_d, p_d => this%p_d, s_hat_d => this%s_hat_d, &
663 p_hat_d => this%p_hat_d, mult_d => coef%mult_d)
665 call device_rzero(x%x_d, n)
666 call device_copy(r_d, f_d, n)
667 rtr = device_glsc3(r_d, mult_d, r_d, n)
673 rnorm = r_norm * norm_fac
674 gamma = rnorm * this%rel_tol
675 ksp_results%res_start = rnorm
676 ksp_results%res_final = rnorm
682 if (r_norm .le. 0.0_rp .or. rnorm .lt. this%abs_tol .or. &
683 rnorm .lt. gamma)
then
684 ksp_results%converged = .true.
692 call this%monitor_start(
'BiCGStab')
693 do iter = 1, max_iter
700 r_norm,
'rho inner product')
702 if (iter .eq. 1)
then
703 call device_copy(p_d, r_d, n)
705 beta = (rho_1 / rho_2) * (alpha / omega)
706 if (.not. ieee_is_finite(beta))
then
707 call neko_error(
'BiCGStab failure: non-finite beta')
712 call this%M%solve(this%p_hat, this%p, n)
713 call ax%compute(this%v, this%p_hat, coef, x%msh, x%Xh)
714 call gs_h%op(this%v, n, gs_op_add, this%gs_event)
715 call device_event_sync(this%gs_event)
716 call bc_projector%apply(this%v, n)
725 v_norm,
'alpha denominator')
727 if (.not. ieee_is_finite(alpha))
then
728 call neko_error(
'BiCGStab failure: non-finite alpha')
734 rnorm = s_norm * norm_fac
735 if (rnorm .lt. this%abs_tol .or. rnorm .lt. gamma)
then
736 call device_add2s2(x%x_d, p_hat_d, alpha, n)
737 call this%monitor_iter(iter, rnorm)
741 call this%M%solve(this%s_hat, this%s, n)
742 call ax%compute(this%t, this%s_hat, coef, x%msh, x%Xh)
743 call gs_h%op(this%t, n, gs_op_add, this%gs_event)
744 call device_event_sync(this%gs_event)
745 call bc_projector%apply(this%t, n)
750 if (t_norm .le. 0.0_rp)
then
751 call neko_error(
'BiCGStab breakdown: zero omega denominator')
753 if (.not. ieee_is_finite(stt))
then
754 call neko_error(
'BiCGStab failure: non-finite omega numerator')
757 if (.not. ieee_is_finite(omega))
then
758 call neko_error(
'BiCGStab failure: non-finite omega')
762 s_hat_d, s_d, t_d, f_d, mult_d, alpha, omega, n)
765 rnorm = r_norm * norm_fac
766 call this%monitor_iter(iter, rnorm)
767 if (rnorm .lt. this%abs_tol .or. rnorm .lt. gamma)
then
781 call this%monitor_stop()
782 ksp_results%res_final = rnorm
783 ksp_results%iter = iter
784 ksp_results%converged = this%is_converged(iter, rnorm)
798 real(kind=rp),
intent(in) :: inner_product
799 real(kind=rp),
intent(in) :: norm_a
800 real(kind=rp),
intent(in) :: norm_b
801 character(len=*),
intent(in) :: quantity
802 real(kind=rp) :: large_norm, small_norm
804 if (.not. ieee_is_finite(inner_product))
then
805 call neko_error(
'BiCGStab failure: non-finite ' // trim(quantity))
808 large_norm =
max(norm_a, norm_b)
809 small_norm = min(norm_a, norm_b)
810 if (large_norm .le. 0.0_rp .or. &
811 abs(inner_product) / large_norm .le. neko_eps * small_norm)
then
812 call neko_error(
'BiCGStab breakdown: near-zero ' // trim(quantity))
822 real(kind=rp),
intent(in) ::
value
823 character(len=*),
intent(in) :: quantity
824 real(kind=rp) :: root
826 if (.not. ieee_is_finite(
value) .or.
value .lt. 0.0_rp)
then
827 call neko_error(
'BiCGStab failure: invalid ' // trim(quantity) // &
852 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
854 class(ax_t),
intent(in) :: ax
855 type(field_t),
intent(inout) :: x
856 type(field_t),
intent(inout) :: y
857 type(field_t),
intent(inout) :: z
858 integer,
intent(in) :: n
859 real(kind=rp),
dimension(n),
intent(in) :: fx
860 real(kind=rp),
dimension(n),
intent(in) :: fy
861 real(kind=rp),
dimension(n),
intent(in) :: fz
862 type(coef_t),
intent(inout) :: coef
863 class(vector_bc_projector_t),
intent(inout) :: bc_projector
864 type(gs_t),
intent(inout) :: gs_h
865 type(ksp_monitor_t),
dimension(3) :: ksp_results
866 integer,
optional,
intent(in) :: niter
867 type(scalar_bc_projector_t),
pointer :: bc_x, bc_y, bc_z
869 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
870 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
871 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
872 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
real_xp opencl_bicgstab_part1(void *s, void *r, void *v, void *mult, real *alpha, int *n, cl_command_queue cmd_queue)
void opencl_bicgstab_part2(void *x, void *r, void *p_hat, void *s_hat, void *s, void *t, void *f, void *mult, real *alpha, real *omega, real_xp *res, int *n, cl_command_queue cmd_queue)
void opencl_bicgstab_product_and_norm(void *a, void *b, void *mult, real_xp *res, int *n, cl_command_queue cmd_queue)
void opencl_bicgstab_update_p(void *p, void *r, void *v, real *beta, real *omega, int *n, cl_command_queue cmd_queue)
void cuda_bicgstab_product_and_norm(void *a, void *b, void *mult, real_xp *res, int *n)
void cuda_bicgstab_update_p(void *p, void *r, void *v, real *beta, real *omega, int *n)
void cuda_bicgstab_part2(void *x, void *r, void *p_hat, void *s_hat, void *s, void *t, void *f, void *mult, real *alpha, real *omega, real_xp *res, int *n)
real_xp cuda_bicgstab_part1(void *s, void *r, void *v, void *mult, real *alpha, int *n)
__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.
Provides a device implementation of the BiCGStab method.
subroutine device_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
Search direction update .
subroutine bicgstab_device_check_inner_product(inner_product, norm_a, norm_b, quantity)
Check an inner product for a BiCGStab breakdown.
type(ksp_monitor_t) function bicgstab_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Solve a linear system with the device BiCGStab method.
subroutine device_bicgstab_product_and_norm(product, norm_squared, a_d, b_d, mult_d, n)
Weighted inner product and squared norm in one reduction.
type(ksp_monitor_t) function, dimension(3) bicgstab_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Solve three independent systems with the device BiCGStab method.
real(kind=rp) function device_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n)
BiCGStab part 1, .
subroutine bicgstab_device_free(this)
Free a device BiCGStab solver.
subroutine bicgstab_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device BiCGStab solver.
real(kind=rp) function bicgstab_device_sqrt(value, quantity)
Return the square root of a valid squared norm.
subroutine device_bicgstab_part2(rtr, rho, x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, mult_d, alpha, omega, n)
BiCGStab part 2, and .
integer, public pe_size
MPI size of communicator.
type(mpi_comm), public neko_comm
MPI communicator.
type(mpi_datatype), public mpi_extra_precision
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.
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
subroutine, public device_event_create(event, flags)
Create a device event queue.
Implements the base abstract type for Krylov solvers plus helper types.
real(kind=rp), parameter, public neko_eps
Machine epsilon .
integer, parameter, public c_xp
integer, parameter, public xp
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 .
Device implementation of the right-preconditioned BiCGStab 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.