52 use,
intrinsic :: iso_c_binding
58 real(kind=
rp),
allocatable :: d(:)
59 real(kind=
rp),
allocatable :: w(:)
60 real(kind=
rp),
allocatable :: r(:)
61 real(kind=
rp) :: omega
64 integer :: max_iter = 10
65 logical :: recompute_diag = .true.
75 real(kind=
rp),
allocatable :: d(:)
76 type(c_ptr) :: d_d = c_null_ptr
77 real(kind=
rp),
allocatable :: w(:)
78 type(c_ptr) :: w_d = c_null_ptr
79 real(kind=
rp),
allocatable :: r(:)
80 type(c_ptr) :: r_d = c_null_ptr
83 real(kind=
rp),
allocatable :: ev(:)
84 type(c_ptr) :: ev_d = c_null_ptr
85 real(kind=
rp) :: tha, dlt
88 integer :: power_its = 250
90 integer :: power_its_refresh = 20
92 logical :: warm_start_eigs = .false.
94 logical :: eigs_computed = .false.
95 integer :: max_iter = 10
96 logical :: recompute_eigs = .true.
115 integer,
intent(in) :: n
116 integer,
intent(in) :: lvl
117 integer,
intent(in) :: max_iter
129 this%max_iter = max_iter
130 this%recompute_eigs = .true.
139 if (
allocated(this%d))
then
145 if (
allocated(this%w))
then
151 if (
allocated(this%r))
then
157 if (
allocated(this%ev))
then
172 integer,
intent(in) :: n
173 real(kind=
rp) :: lam, b, a, rn
174 real(kind=
rp),
parameter :: boost = 1.1_rp
175 real(kind=
rp),
parameter :: lam_factor = 30.0_rp
176 real(kind=
rp) :: wtw, dtw, dtd
177 integer,
allocatable :: fixed_seed(:), saved_seed(:)
178 integer :: i, rnd_n, its
182 warm = this%warm_start_eigs .and. this%eigs_computed .and. &
184 associate(w => this%w, d => this%d, coef => amg%coef, gs_h => amg%gs_h, &
185 msh => amg%msh, xh => amg%Xh, blst => amg%blst)
188 its = this%power_its_refresh
189 call copy(d, this%ev, n)
194 call random_seed(
size = rnd_n)
195 allocate(saved_seed(rnd_n))
196 allocate(fixed_seed(rnd_n))
198 call random_seed(get = saved_seed)
199 call random_seed(put = fixed_seed)
202 call random_number(rn)
207 call random_seed(put = saved_seed)
209 if (this%lvl .eq. 0)
then
210 call gs_h%op(d, n, gs_op_add)
211 call blst%apply(d, n)
216 call amg%matvec(w, d, this%lvl)
218 if (this%lvl .eq. 0)
then
219 wtw =
glsc3(w, coef%mult, w, n)
224 call cmult2(d, w, 1.0_rp/sqrt(wtw), n)
227 call amg%matvec(w, d, this%lvl)
229 if (this%lvl .eq. 0)
then
230 dtw =
glsc3(d, coef%mult, w, n)
231 dtd =
glsc3(d, coef%mult, d, n)
239 this%tha = (b+a)/2.0_rp
240 this%dlt = (b-a)/2.0_rp
242 if (this%warm_start_eigs)
then
243 if (.not.
allocated(this%ev))
then
244 allocate(this%ev(this%n))
246 call copy(this%ev, d, n)
248 this%eigs_computed = .true.
250 this%recompute_eigs = .false.
264 integer,
intent(in) :: n
265 real(kind=rp),
dimension(n),
intent(inout) :: x
266 real(kind=rp),
dimension(n),
intent(inout) :: f
267 class(tamg_hierarchy_t),
intent(inout) :: amg
268 type(ksp_monitor_t) :: ksp_results
269 logical,
optional,
intent(in) :: zero_init
270 integer :: iter, max_iter, i
271 real(kind=rp) :: rtr, rnorm
272 real(kind=rp) :: rhok, rhokp1, s1, thet, delt, tmp1, tmp2
273 logical :: zero_initial_guess
275 if (this%recompute_eigs)
then
276 call this%comp_eig(amg, n)
278 if (
present(zero_init))
then
279 zero_initial_guess = zero_init
281 zero_initial_guess = .false.
283 max_iter = this%max_iter
285 associate( w => this%w, r => this%r, d => this%d, blst => amg%blst)
287 if (.not. zero_initial_guess)
then
288 call amg%matvec(w, x, this%lvl)
304 d(i) = 1.0_rp/thet * r(i)
310 do iter = 2, max_iter
311 call amg%matvec(w, d, this%lvl)
313 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
315 tmp2 = 2.0_rp * rhokp1 / delt
326 d(i) = tmp1 * d(i) + tmp2 * r(i)
340 type(tamg_hierarchy_t),
intent(inout) :: amg
341 integer,
intent(in) :: n
342 real(kind=rp) :: lam, b, a, rn
343 real(kind=rp),
parameter :: boost = 1.1_rp
344 real(kind=rp),
parameter :: lam_factor = 30.0_rp
345 real(kind=rp) :: wtw, dtw, dtd
346 integer,
allocatable :: fixed_seed(:), saved_seed(:)
347 integer :: i, rnd_n, its
350 call profiler_start_region(
'AMG_cheby_power')
351 warm = this%warm_start_eigs .and. this%eigs_computed .and. &
352 c_associated(this%ev_d)
353 associate(w => this%w, d => this%d, coef => amg%coef, gs_h => amg%gs_h, &
354 msh => amg%msh, xh => amg%Xh, blst => amg%blst)
357 its = this%power_its_refresh
358 call device_copy(this%d_d, this%ev_d, n)
363 call random_seed(
size = rnd_n)
364 allocate(saved_seed(rnd_n))
365 allocate(fixed_seed(rnd_n))
367 call random_seed(get = saved_seed)
368 call random_seed(put = fixed_seed)
371 call random_number(rn)
374 call device_memcpy(this%d, this%d_d, n, host_to_device, .true.)
377 call random_seed(put = saved_seed)
379 if (this%lvl .eq. 0)
then
380 call gs_h%op(d, n, gs_op_add)
381 call blst%apply(d, n)
385 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
387 if (this%lvl .eq. 0)
then
388 wtw = device_glsc3(this%w_d, coef%mult_d, this%w_d, n)
390 wtw = device_glsc2(this%w_d, this%w_d, n)
393 call device_cmult2(this%d_d, this%w_d, 1.0_rp/sqrt(wtw), n)
396 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
398 if (this%lvl .eq. 0)
then
399 dtw = device_glsc3(this%d_d, coef%mult_d, this%w_d, n)
400 dtd = device_glsc3(this%d_d, coef%mult_d, this%d_d, n)
402 dtw = device_glsc2(this%d_d, this%w_d, n)
403 dtd = device_glsc2(this%d_d, this%d_d, n)
408 this%tha = (b+a)/2.0_rp
409 this%dlt = (b-a)/2.0_rp
411 if (this%warm_start_eigs)
then
412 if (.not. c_associated(this%ev_d))
then
413 allocate(this%ev(this%n))
414 call device_map(this%ev, this%ev_d, this%n)
416 call device_copy(this%ev_d, this%d_d, n)
418 this%eigs_computed = .true.
420 this%recompute_eigs = .false.
423 call profiler_end_region(
'AMG_cheby_power')
434 integer,
intent(in) :: n
435 real(kind=rp),
dimension(n),
intent(inout) :: x
436 real(kind=rp),
dimension(n),
intent(inout) :: f
439 class(tamg_hierarchy_t),
intent(inout) :: amg
440 type(ksp_monitor_t) :: ksp_results
441 logical,
optional,
intent(in) :: zero_init
442 integer :: iter, max_iter
443 real(kind=rp) :: rtr, rnorm
444 real(kind=rp) :: rhok, rhokp1, s1, thet, delt, tmp1, tmp2
445 logical :: zero_initial_guess
447 if (this%recompute_eigs)
then
448 call this%device_comp_eig(amg, n)
450 if (
present(zero_init))
then
451 zero_initial_guess = zero_init
453 zero_initial_guess = .false.
455 max_iter = this%max_iter
457 associate( w_d => this%w_d, r_d => this%r_d, d_d => this%d_d, &
460 if (.not. zero_initial_guess)
then
461 call amg%device_matvec(this%w, x, w_d, x_d, this%lvl)
471 call amg_device_cheby_solve_part1(r_d, f_d, w_d, x_d, d_d, &
472 tmp1, n, zero_initial_guess)
474 do iter = 2, max_iter
475 call amg%device_matvec(this%w, this%d, w_d, d_d, this%lvl)
477 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
479 tmp2 = 2.0_rp * rhokp1 / delt
482 call amg_device_cheby_solve_part2(r_d, w_d, d_d, x_d, tmp1, tmp2, n)
495 integer,
intent(in) :: n
496 integer,
intent(in) :: lvl
497 integer,
intent(in) :: max_iter
504 this%max_iter = max_iter
512 if (
allocated(this%d))
then
515 if (
allocated(this%w))
then
518 if (
allocated(this%r))
then
528 type(tamg_hierarchy_t),
intent(inout) :: amg
529 integer,
intent(in) :: n
533 call tamg_sample_matrix_val(val, amg, this%lvl, i, i)
534 this%d(i) = 1.0_rp / val
536 this%recompute_diag = .false.
546 integer,
intent(in) :: n
547 real(kind=rp),
dimension(n),
intent(inout) :: x
548 real(kind=rp),
dimension(n),
intent(inout) :: f
549 class(tamg_hierarchy_t),
intent(inout) :: amg
550 type(ksp_monitor_t) :: ksp_results
551 integer,
optional,
intent(in) :: niter
552 integer :: iter, max_iter
553 real(kind=rp) :: rtr, rnorm
556 if (this%recompute_diag)
then
557 call this%comp_diag(amg, n)
560 if (
present(niter))
then
563 max_iter = this%max_iter
567 associate( w => this%w, r => this%r, d => this%d)
568 do iter = 1, max_iter
571 call amg%matvec(w, x, this%lvl)
588 x(i) = x(i) + this%omega * r(i)
597 integer,
intent(in) :: lvl
599 character(len=LOG_SIZE) :: log_buf
601 write(log_buf,
'(A8,I2,A28)')
'-- level', lvl,
'-- init smoother: Chebyshev'
602 call neko_log%message(log_buf)
603 write(log_buf,
'(A22,I6)')
'Iterations:', smoo%max_iter
604 call neko_log%message(log_buf)
608 integer,
intent(in) :: lvl
609 real(kind=rp),
intent(in) :: lam
610 character(len=LOG_SIZE) :: log_buf
612 write(log_buf,
'(A12,I2,A29,F12.3)')
'-- AMG level', lvl, &
613 '-- Chebyshev approx. max eig', lam
614 call neko_log%message(log_buf)
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
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)
subroutine, public device_add2(a_d, b_d, n, strm)
Vector addition .
subroutine, public device_add3s2(a_d, b_d, c_d, c1, c2, n, strm)
Returns .
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
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 .
real(kind=rp) function, public device_glsc2(a_d, b_d, n, strm)
Weighted inner product .
subroutine, public device_cmult2(a_d, b_d, c, n, strm)
Multiplication by constant c .
Implements device kernels for use with TreeAMG smoothers.
subroutine, public amg_device_cheby_solve_part1(r_d, f_d, w_d, x_d, d_d, inv_thet, n, zero_initial)
subroutine, public amg_device_cheby_solve_part2(r_d, w_d, d_d, x_d, tmp1, tmp2, n)
Device abstraction, common interface for various accelerators.
integer, parameter, public host_to_device
Implements the base abstract type for Krylov solvers plus helper types.
type(log_t), public neko_log
Global log stream.
integer, parameter, public log_size
subroutine, public cmult(a, c, n)
Multiplication by constant c .
subroutine, public cmult2(a, b, c, n)
Multiplication by constant c .
real(kind=rp) function, public glsc3(a, b, c, n)
Weighted inner product .
real(kind=rp) function, public glsc2(a, b, n)
Weighted inner product .
subroutine, public add2(a, b, n)
Vector addition .
subroutine, public add3s2(a, b, c, c1, c2, n)
Returns .
subroutine, public col2(a, b, n)
Vector multiplication .
subroutine, public copy(a, b, n)
Copy a vector .
subroutine, public sub2(a, b, n)
Vector substraction .
subroutine, public add2s2(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on second argument)
integer, parameter neko_bcknd_device
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 smoothers for use with TreeAMG matrix vector product.
subroutine amg_cheby_free(this)
free cheby data
subroutine amg_jacobi_init(this, n, lvl, max_iter)
Initialization of Jacobi (this is expensive...)
subroutine amg_cheby_power(this, amg, n)
Power method to approximate largest eigenvalue.
subroutine amg_device_cheby_power(this, amg, n)
Power method to approximate largest eigenvalue.
subroutine amg_cheby_solve(this, x, f, n, amg, zero_init)
Chebyshev smoother From Saad's iterative methods textbook.
subroutine amg_cheby_init(this, n, lvl, max_iter)
Initialization of chebyshev.
subroutine amg_jacobi_free(this)
free jacobi data
subroutine amg_jacobi_diag(this, amg, n)
SAMPLE MATRIX DIAGONAL VALUES (DO NOT USE, EXPENSIVE)
subroutine amg_device_cheby_solve(this, x, f, x_d, f_d, n, amg, zero_init)
Chebyshev smoother From Saad's iterative methods textbook.
subroutine amg_jacobi_solve(this, x, f, n, amg, niter)
Jacobi smoother.
subroutine amg_smoo_monitor(lvl, smoo)
subroutine amg_cheby_monitor(lvl, lam)
Implements utilities for the TreeAMG hierarchy structure.
subroutine, public tamg_sample_matrix_val(val, amg, lvl, i, j)
Sample the values in a matix (expensive, use with caution)
Implements the base type for TreeAMG hierarchy structure.
A list of allocatable `bc_t`. Follows the standard interface of lists.
Type for storing initial and final residuals in a Krylov solver.
Type for a TreeAMG hierarchy.
Type for Chebyshev iteration using TreeAMG matvec.
Type for Chebyshev iteration using TreeAMG matvec.