55 use,
intrinsic :: iso_c_binding, only : c_ptr, c_int, &
56 c_null_ptr, c_associated
62 real(kind=
rp),
allocatable :: d(:)
63 real(kind=
rp),
allocatable :: w(:)
64 real(kind=
rp),
allocatable :: r(:)
65 type(c_ptr) :: d_d = c_null_ptr
66 type(c_ptr) :: w_d = c_null_ptr
67 type(c_ptr) :: r_d = c_null_ptr
68 type(c_ptr) :: gs_event = c_null_ptr
69 real(kind=
rp) :: tha, dlt
70 integer :: power_its = 150
72 integer :: power_its_refresh = 20
74 logical :: warm_start_eigs = .false.
76 logical :: eigs_computed = .false.
78 real(kind=
rp),
allocatable :: ev(:)
79 type(c_ptr) :: ev_d = c_null_ptr
80 logical :: recompute_eigs = .true.
81 logical :: zero_initial_guess = .false.
93 bind(c, name =
'hip_cheby_part1')
94 use,
intrinsic :: iso_c_binding
97 type(c_ptr),
value :: d_d, x_d, strm
105 bind(c, name =
'hip_cheby_part2')
106 use,
intrinsic :: iso_c_binding
109 type(c_ptr),
value :: d_d, w_d, x_d, strm
110 real(c_rp) :: tmp1, tmp2
116 subroutine cuda_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
117 bind(c, name =
'cuda_cheby_part1')
118 use,
intrinsic :: iso_c_binding
121 type(c_ptr),
value :: d_d, x_d, strm
122 real(c_rp) :: inv_tha
124 end subroutine cuda_cheby_device_part1
128 subroutine cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
129 bind(c, name =
'cuda_cheby_part2')
130 use,
intrinsic :: iso_c_binding
133 type(c_ptr),
value :: d_d, w_d, x_d, strm
134 real(c_rp) :: tmp1, tmp2
136 end subroutine cuda_cheby_device_part2
140 subroutine metal_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
141 bind(c, name =
'metal_cheby_part1')
142 use,
intrinsic :: iso_c_binding
145 type(c_ptr),
value :: d_d, x_d, strm
146 real(c_rp) :: inv_tha
148 end subroutine metal_cheby_device_part1
152 subroutine metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
153 bind(c, name =
'metal_cheby_part2')
154 use,
intrinsic :: iso_c_binding
157 type(c_ptr),
value :: d_d, w_d, x_d, strm
158 real(c_rp) :: tmp1, tmp2
160 end subroutine metal_cheby_device_part2
166 type(c_ptr) :: d_d, x_d
167 real(c_rp) :: inv_tha
172 call cuda_cheby_device_part1(d_d, x_d, inv_tha, n,
glb_cmd_queue)
174 call metal_cheby_device_part1(d_d, x_d, inv_tha, n,
glb_cmd_queue)
175#else !Fallback to device_math for missing device kernels
184 type(c_ptr) :: d_d, w_d, x_d
185 real(c_rp) :: tmp1, tmp2
190 call cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n,
glb_cmd_queue)
192 call metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n,
glb_cmd_queue)
193#else !Fallback to device_math for missing device kernels
203 integer,
intent(in) :: max_iter
204 class(
pc_t),
optional,
intent(in),
target :: M
205 integer,
intent(in) :: n
206 real(kind=
rp),
optional,
intent(in) :: rel_tol
207 real(kind=
rp),
optional,
intent(in) :: abs_tol
208 logical,
optional,
intent(in) :: monitor
223 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
224 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
225 else if (
present(rel_tol) .and.
present(abs_tol))
then
226 call this%ksp_init(max_iter, rel_tol, abs_tol)
227 else if (
present(monitor) .and.
present(abs_tol))
then
228 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
229 else if (
present(rel_tol) .and.
present(monitor))
then
230 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
231 else if (
present(rel_tol))
then
232 call this%ksp_init(max_iter, rel_tol = rel_tol)
233 else if (
present(abs_tol))
then
234 call this%ksp_init(max_iter, abs_tol = abs_tol)
235 else if (
present(monitor))
then
236 call this%ksp_init(max_iter, monitor = monitor)
238 call this%ksp_init(max_iter)
250 if (
allocated(this%d))
then
251 if (c_associated(this%d_d))
then
257 if (
allocated(this%w))
then
258 if (c_associated(this%w_d))
then
264 if (
allocated(this%r))
then
265 if (c_associated(this%r_d))
then
271 if (
allocated(this%ev))
then
272 if (c_associated(this%ev_d))
then
280 if (c_associated(this%gs_event))
then
288 class(
ax_t),
intent(in) :: Ax
289 type(
field_t),
intent(inout) :: x
290 integer,
intent(in) :: n
291 type(
coef_t),
intent(inout) :: coef
293 type(
gs_t),
intent(inout) :: gs_h
294 real(kind=
rp) :: lam, b, a, rn
295 real(kind=
rp) :: boost = 1.1_rp
296 real(kind=
rp) :: lam_factor = 30.0_rp
297 real(kind=
rp) :: wtw, dtw, dtd
298 integer,
allocatable :: fixed_seed(:), saved_seed(:)
299 integer :: i, rnd_n, its
303 warm = this%warm_start_eigs .and. this%eigs_computed .and. &
304 c_associated(this%ev_d)
305 associate(w => this%w, w_d => this%w_d, d => this%d, d_d => this%d_d)
308 its = this%power_its_refresh
314 call random_seed(
size = rnd_n)
315 allocate(saved_seed(rnd_n))
316 allocate(fixed_seed(rnd_n))
318 call random_seed(get = saved_seed)
319 call random_seed(put = fixed_seed)
322 call random_number(rn)
328 call random_seed(put = saved_seed)
330 call gs_h%op(d, n, gs_op_add, this%gs_event)
331 call bc_projector%apply(d, n)
336 call ax%compute(w, d, coef, x%msh, x%Xh)
337 call gs_h%op(w, n, gs_op_add, this%gs_event)
338 call bc_projector%apply(w, n)
339 if (
associated(this%schwarz))
then
340 call this%schwarz%compute(this%r, w)
343 call this%M%solve(this%r, w, n)
349 call bc_projector%apply(d, n)
352 call ax%compute(w, d, coef, x%msh, x%Xh)
353 call gs_h%op(w, n, gs_op_add, this%gs_event)
354 call bc_projector%apply(w, n)
355 if (
associated(this%schwarz))
then
356 call this%schwarz%compute(this%r, w)
359 call this%M%solve(this%r, w, n)
368 this%tha = (b+a)/2.0_rp
369 this%dlt = (b-a)/2.0_rp
371 if (this%warm_start_eigs)
then
372 if (.not. c_associated(this%ev_d))
then
373 allocate(this%ev(
size(this%d)))
374 call device_map(this%ev, this%ev_d,
size(this%d))
378 this%eigs_computed = .true.
380 this%recompute_eigs = .false.
387 niter)
result(ksp_results)
389 class(ax_t),
intent(in) :: ax
390 type(field_t),
intent(inout) :: x
391 integer,
intent(in) :: n
392 real(kind=rp),
dimension(n),
intent(in) :: f
393 type(coef_t),
intent(inout) :: coef
394 class(scalar_bc_projector_t),
intent(inout) :: bc_projector
395 type(gs_t),
intent(inout) :: gs_h
396 type(ksp_monitor_t) :: ksp_results
397 integer,
optional,
intent(in) :: niter
398 integer :: iter, max_iter
399 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
402 f_d = device_get_ptr(f)
404 if (this%recompute_eigs)
then
408 if (
present(niter))
then
411 max_iter = this%max_iter
413 norm_fac = 1.0_rp / sqrt(coef%volume)
415 associate( w => this%w, r => this%r, d => this%d, &
416 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
418 call device_copy(r_d, f_d, n)
419 call ax%compute(w, x%x, coef, x%msh, x%Xh)
420 call gs_h%op(w, n, gs_op_add, this%gs_event)
421 call bc_projector%apply(w, n)
422 call device_sub2(r_d, w_d, n)
424 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
425 rnorm = sqrt(rtr) * norm_fac
426 ksp_results%res_start = rnorm
427 ksp_results%res_final = rnorm
431 call this%M%solve(w, r, n)
432 call device_copy(d_d, w_d, n)
433 a = 2.0_rp / this%tha
434 call device_add2s2(x%x_d, d_d, a, n)
437 do iter = 2, max_iter
439 call device_copy(r_d, f_d, n)
440 call ax%compute(w, x%x, coef, x%msh, x%Xh)
441 call gs_h%op(w, n, gs_op_add, this%gs_event)
442 call bc_projector%apply(w, n)
443 call device_sub2(r_d, w_d, n)
445 call this%M%solve(w, r, n)
447 if (iter .eq. 2)
then
448 b = 0.5_rp * (this%dlt * a)**2
450 b = (this%dlt * a / 2.0_rp)**2
452 a = 1.0_rp/(this%tha - b/a)
453 call device_add2s1(d_d, w_d, b, n)
455 call device_add2s2(x%x_d, d_d, a, n)
459 call device_copy(r_d, f_d, n)
460 call ax%compute(w, x%x, coef, x%msh, x%Xh)
461 call gs_h%op(w, n, gs_op_add, this%gs_event)
462 call bc_projector%apply(w, n)
463 call device_sub2(r_d, w_d, n)
464 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
465 rnorm = sqrt(rtr) * norm_fac
468 ksp_results%res_final = rnorm
469 ksp_results%iter = iter
470 ksp_results%converged = this%is_converged(iter, rnorm)
476 niter)
result(ksp_results)
478 class(ax_t),
intent(in) :: ax
479 type(field_t),
intent(inout) :: x
480 integer,
intent(in) :: n
481 real(kind=rp),
dimension(n),
intent(in) :: f
482 type(coef_t),
intent(inout) :: coef
483 class(scalar_bc_projector_t),
intent(inout) :: bc_projector
484 type(gs_t),
intent(inout) :: gs_h
485 type(ksp_monitor_t) :: ksp_results
486 integer,
optional,
intent(in) :: niter
487 integer :: iter, max_iter
488 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
489 real(kind=rp) :: rhok, rhokp1, sig1, tmp1, tmp2
492 f_d = device_get_ptr(f)
494 if (this%recompute_eigs)
then
498 if (
present(niter))
then
501 max_iter = this%max_iter
503 norm_fac = 1.0_rp / sqrt(coef%volume)
505 associate( w => this%w, r => this%r, d => this%d, &
506 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
508 if (.not.this%zero_initial_guess)
then
509 call ax%compute(w, x%x, coef, x%msh, x%Xh)
510 call gs_h%op(w, n, gs_op_add, this%gs_event)
511 call bc_projector%apply(w, n)
512 call device_sub3(r_d, f_d, w_d, n)
514 call device_copy(r_d, f_d, n)
515 this%zero_initial_guess = .false.
519 if (
associated(this%schwarz))
then
520 call this%schwarz%compute(d, r)
522 call this%M%solve(d, r, n)
525 tmp1 = 1.0_rp / this%tha
528 sig1 = this%tha / this%dlt
532 do iter = 2, max_iter
533 rhokp1 = 1.0_rp / (2.0_rp * sig1 - rhok)
535 tmp2 = 2.0_rp * rhokp1 / this%dlt
538 call ax%compute(w, x%x, coef, x%msh, x%Xh)
539 call gs_h%op(w, n, gs_op_add, this%gs_event)
540 call bc_projector%apply(w, n)
541 call device_sub3(r_d, f_d, w_d, n)
543 if (
associated(this%schwarz))
then
544 call this%schwarz%compute(w, r)
546 call this%M%solve(w, r, n)
558 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
560 class(ax_t),
intent(in) :: ax
561 type(field_t),
intent(inout) :: x
562 type(field_t),
intent(inout) :: y
563 type(field_t),
intent(inout) :: z
564 integer,
intent(in) :: n
565 real(kind=rp),
dimension(n),
intent(in) :: fx
566 real(kind=rp),
dimension(n),
intent(in) :: fy
567 real(kind=rp),
dimension(n),
intent(in) :: fz
568 type(coef_t),
intent(inout) :: coef
569 class(vector_bc_projector_t),
intent(inout) :: bc_projector
570 type(gs_t),
intent(inout) :: gs_h
571 type(ksp_monitor_t),
dimension(3) :: ksp_results
572 integer,
optional,
intent(in) :: niter
573 type(scalar_bc_projector_t),
pointer :: bc_x, bc_y, bc_z
575 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
576 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
577 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
578 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.
Chebyshev preconditioner.
subroutine cheby_device_power(this, ax, x, n, coef, bc_projector, gs_h)
type(ksp_monitor_t) function cheby_device_impl(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
A chebyshev preconditioner.
subroutine cheby_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a standard solver.
subroutine cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n)
subroutine cheby_device_free(this)
subroutine cheby_device_part1(d_d, x_d, inv_tha, n)
type(ksp_monitor_t) function cheby_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
A chebyshev preconditioner.
type(ksp_monitor_t) function, dimension(3) cheby_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard Cheby_Deviceshev coupled solve.
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_sub3(a_d, b_d, c_d, n, strm)
Vector subtraction .
subroutine, public device_add2s2(a_d, b_d, c1, n, strm)
Vector addition with scalar multiplication (multiplication on first argument)
subroutine, public device_add2(a_d, b_d, n, strm)
Vector addition .
subroutine, public device_cmult(a_d, c, n, strm)
Multiplication by constant c .
subroutine, public device_sub2(a_d, b_d, n, strm)
Vector substraction .
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 .
subroutine, public device_cmult2(a_d, b_d, c, n, strm)
Multiplication by constant c .
Device abstraction, common interface for various accelerators.
integer, parameter, public host_to_device
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.
integer, parameter, public c_rp
integer, parameter, public rp
Global precision used in computations.
subroutine, public profiler_start_region(name, region_id)
Started a named (name) profiler region.
subroutine, public profiler_end_region(name, region_id)
End the most recently started profiler region.
Implements scalar_projector_t.
Overlapping schwarz solves.
Defines a function space.
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 .
Defines a Chebyshev preconditioner.
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.
The function space for the SEM solution fields.
Abstract type for resolving vector boundary conditions.