Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
tree_amg_smoother.f90
Go to the documentation of this file.
1! Copyright (c) 2024-2026, The Neko Authors
2! All rights reserved.
3!
4! Redistribution and use in source and binary forms, with or without
5! modification, are permitted provided that the following conditions
6! are met:
7!
8! * Redistributions of source code must retain the above copyright
9! notice, this list of conditions and the following disclaimer.
10!
11! * Redistributions in binary form must reproduce the above
12! copyright notice, this list of conditions and the following
13! disclaimer in the documentation and/or other materials provided
14! with the distribution.
15!
16! * Neither the name of the authors nor the names of its
17! contributors may be used to endorse or promote products derived
18! from this software without specific prior written permission.
19!
20! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
21! "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
22! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
23! FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
24! COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
25! INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
26! BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
27! LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
28! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
29! LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
30! ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
31! POSSIBILITY OF SUCH DAMAGE.
32!
35 use tree_amg, only : tamg_hierarchy_t
37 use num_types, only : rp
38 use math, only : col2, add2, add2s2, glsc2, glsc3, sub2, cmult, &
43 use krylov, only : ksp_monitor_t
44 use bc_list, only: bc_list_t
45 use gather_scatter, only : gs_t, gs_op_add
46 use logger, only : neko_log, log_size
52 use, intrinsic :: iso_c_binding
53 implicit none
54 private
55
57 type, public :: amg_jacobi_t
58 real(kind=rp), allocatable :: d(:)
59 real(kind=rp), allocatable :: w(:)
60 real(kind=rp), allocatable :: r(:)
61 real(kind=rp) :: omega
62 integer :: lvl
63 integer :: n
64 integer :: max_iter = 10
65 logical :: recompute_diag = .true.
66 contains
67 procedure, pass(this) :: init => amg_jacobi_init
68 procedure, pass(this) :: solve => amg_jacobi_solve
69 procedure, pass(this) :: comp_diag => amg_jacobi_diag
70 procedure, pass(this) :: free => amg_jacobi_free
71 end type amg_jacobi_t
72
74 type, public :: amg_cheby_t
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
86 integer :: lvl
87 integer :: n
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.
97 contains
98 procedure, pass(this) :: init => amg_cheby_init
99 procedure, pass(this) :: solve => amg_cheby_solve
100 procedure, pass(this) :: comp_eig => amg_cheby_power
101 procedure, pass(this) :: device_solve => amg_device_cheby_solve
102 procedure, pass(this) :: device_comp_eig => amg_device_cheby_power
103 procedure, pass(this) :: free => amg_cheby_free
104 end type amg_cheby_t
105
106contains
107
113 subroutine amg_cheby_init(this, n, lvl, max_iter)
114 class(amg_cheby_t), intent(inout), target :: this
115 integer, intent(in) :: n
116 integer, intent(in) :: lvl
117 integer, intent(in) :: max_iter
118
119 allocate(this%d(n))
120 allocate(this%w(n))
121 allocate(this%r(n))
122 if (neko_bcknd_device .eq. 1) then
123 call device_map(this%d, this%d_d, n)
124 call device_map(this%w, this%w_d, n)
125 call device_map(this%r, this%r_d, n)
126 end if
127 this%n = n
128 this%lvl = lvl
129 this%max_iter = max_iter
130 this%recompute_eigs = .true.
131
132 call amg_smoo_monitor(lvl, this)
133
134 end subroutine amg_cheby_init
135
137 subroutine amg_cheby_free(this)
138 class(amg_cheby_t), intent(inout), target :: this
139 if (allocated(this%d)) then
140 if (neko_bcknd_device .eq. 1 .and. c_associated(this%d_d)) then
141 call device_unmap(this%d, this%d_d)
142 end if
143 deallocate(this%d)
144 end if
145 if (allocated(this%w)) then
146 if (neko_bcknd_device .eq. 1 .and. c_associated(this%w_d)) then
147 call device_unmap(this%w, this%w_d)
148 end if
149 deallocate(this%w)
150 end if
151 if (allocated(this%r)) then
152 if (neko_bcknd_device .eq. 1 .and. c_associated(this%r_d)) then
153 call device_unmap(this%r, this%r_d)
154 end if
155 deallocate(this%r)
156 end if
157 if (allocated(this%ev)) then
158 if (neko_bcknd_device .eq. 1 .and. c_associated(this%ev_d)) then
159 call device_unmap(this%ev, this%ev_d)
160 end if
161 deallocate(this%ev)
162 end if
163 end subroutine amg_cheby_free
164
165
169 subroutine amg_cheby_power(this, amg, n)
170 class(amg_cheby_t), intent(inout) :: this
171 type(tamg_hierarchy_t), intent(inout) :: amg
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
179 logical :: warm
180
181 call profiler_start_region('AMG_cheby_power')
182 warm = this%warm_start_eigs .and. this%eigs_computed .and. &
183 allocated(this%ev)
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)
186
187 if (warm) then
188 its = this%power_its_refresh
189 call copy(d, this%ev, n)
190 else
191 its = this%power_its
192
193 ! Save current random seed and set a fixed seed
194 call random_seed(size = rnd_n)
195 allocate(saved_seed(rnd_n))
196 allocate(fixed_seed(rnd_n))
197 fixed_seed = 3901
198 call random_seed(get = saved_seed)
199 call random_seed(put = fixed_seed)
200
201 do i = 1, n
202 call random_number(rn)
203 d(i) = rn + 10.0_rp
204 end do
205
206 ! Restore saved random seed
207 call random_seed(put = saved_seed)
208
209 if (this%lvl .eq. 0) then
210 call gs_h%op(d, n, gs_op_add)!TODO
211 call blst%apply(d, n)
212 end if
213 end if
214 !Power method to get lamba max
215 do i = 1, its
216 call amg%matvec(w, d, this%lvl)
217
218 if (this%lvl .eq. 0) then
219 wtw = glsc3(w, coef%mult, w, n)
220 else
221 wtw = glsc2(w, w, n)
222 end if
223
224 call cmult2(d, w, 1.0_rp/sqrt(wtw), n)
225 end do
226
227 call amg%matvec(w, d, this%lvl)
228
229 if (this%lvl .eq. 0) then
230 dtw = glsc3(d, coef%mult, w, n)
231 dtd = glsc3(d, coef%mult, d, n)
232 else
233 dtw = glsc2(d, w, n)
234 dtd = glsc2(d, d, n)
235 end if
236 lam = dtw / dtd
237 b = lam * boost
238 a = lam / lam_factor
239 this%tha = (b+a)/2.0_rp
240 this%dlt = (b-a)/2.0_rp
241
242 if (this%warm_start_eigs) then
243 if (.not. allocated(this%ev)) then
244 allocate(this%ev(this%n))
245 end if
246 call copy(this%ev, d, n)
247 end if
248 this%eigs_computed = .true.
249
250 this%recompute_eigs = .false.
251 call amg_cheby_monitor(this%lvl, lam)
252 end associate
253 call profiler_end_region('AMG_cheby_power')
254 end subroutine amg_cheby_power
255
262 subroutine amg_cheby_solve(this, x, f, n, amg, zero_init)
263 class(amg_cheby_t), intent(inout) :: this
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
274
275 if (this%recompute_eigs) then
276 call this%comp_eig(amg, n)
277 end if
278 if (present(zero_init)) then
279 zero_initial_guess = zero_init
280 else
281 zero_initial_guess = .false.
282 end if
283 max_iter = this%max_iter
284
285 associate( w => this%w, r => this%r, d => this%d, blst => amg%blst)
286 call copy(r, f, n)
287 if (.not. zero_initial_guess) then
288 call amg%matvec(w, x, this%lvl)
289 call sub2(r, w, n)
290 end if
291
292 thet = this%tha
293 delt = this%dlt
294 s1 = thet / delt
295 rhok = 1.0_rp / s1
296
297 ! First iteration
298 !OCL NORECURRENCE, NOVREC, NOALIAS
299 !DIR$ CONCURRENT
300 !DIR$ IVDEP
301 !GCC$ ivdep
302 !$omp parallel do
303 do i = 1, n
304 d(i) = 1.0_rp/thet * r(i)
305 x(i) = x(i) + d(i)
306 end do
307 !$omp end parallel do
308
309 ! Rest of iterations
310 do iter = 2, max_iter
311 call amg%matvec(w, d, this%lvl)
312
313 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
314 tmp1 = rhokp1 * rhok
315 tmp2 = 2.0_rp * rhokp1 / delt
316 rhok = rhokp1
317
318 !$omp parallel private(i)
319 !OCL NORECURRENCE, NOVREC, NOALIAS
320 !DIR$ CONCURRENT
321 !DIR$ IVDEP
322 !GCC$ ivdep
323 !$omp do
324 do i = 1, n
325 r(i) = r(i) - w(i)
326 d(i) = tmp1 * d(i) + tmp2 * r(i)
327 x(i) = x(i) + d(i)
328 end do
329 !$omp end do
330 !$omp end parallel
331 end do
332 end associate
333 end subroutine amg_cheby_solve
334
338 subroutine amg_device_cheby_power(this, amg, n)
339 class(amg_cheby_t), intent(inout) :: this
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
348 logical :: warm
349
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)
355
356 if (warm) then
357 its = this%power_its_refresh
358 call device_copy(this%d_d, this%ev_d, n)
359 else
360 its = this%power_its
361
362 ! Save current random seed and set a fixed seed
363 call random_seed(size = rnd_n)
364 allocate(saved_seed(rnd_n))
365 allocate(fixed_seed(rnd_n))
366 fixed_seed = 3901
367 call random_seed(get = saved_seed)
368 call random_seed(put = fixed_seed)
369
370 do i = 1, n
371 call random_number(rn)
372 d(i) = rn + 10.0_rp
373 end do
374 call device_memcpy(this%d, this%d_d, n, host_to_device, .true.)
375
376 ! Restore saved random seed
377 call random_seed(put = saved_seed)
378
379 if (this%lvl .eq. 0) then
380 call gs_h%op(d, n, gs_op_add)!TODO
381 call blst%apply(d, n)
382 end if
383 end if
384 do i = 1, its
385 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
386
387 if (this%lvl .eq. 0) then
388 wtw = device_glsc3(this%w_d, coef%mult_d, this%w_d, n)
389 else
390 wtw = device_glsc2(this%w_d, this%w_d, n)
391 end if
392
393 call device_cmult2(this%d_d, this%w_d, 1.0_rp/sqrt(wtw), n)
394 end do
395
396 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
397
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)
401 else
402 dtw = device_glsc2(this%d_d, this%w_d, n)
403 dtd = device_glsc2(this%d_d, this%d_d, n)
404 end if
405 lam = dtw / dtd
406 b = lam * boost
407 a = lam / lam_factor
408 this%tha = (b+a)/2.0_rp
409 this%dlt = (b-a)/2.0_rp
410
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)
415 end if
416 call device_copy(this%ev_d, this%d_d, n)
417 end if
418 this%eigs_computed = .true.
419
420 this%recompute_eigs = .false.
421 call amg_cheby_monitor(this%lvl, lam)
422 end associate
423 call profiler_end_region('AMG_cheby_power')
424 end subroutine amg_device_cheby_power
425
432 subroutine amg_device_cheby_solve(this, x, f, x_d, f_d, n, amg, zero_init)
433 class(amg_cheby_t), intent(inout) :: this
434 integer, intent(in) :: n
435 real(kind=rp), dimension(n), intent(inout) :: x
436 real(kind=rp), dimension(n), intent(inout) :: f
437 type(c_ptr) :: x_d
438 type(c_ptr) :: f_d
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
446
447 if (this%recompute_eigs) then
448 call this%device_comp_eig(amg, n)
449 end if
450 if (present(zero_init)) then
451 zero_initial_guess = zero_init
452 else
453 zero_initial_guess = .false.
454 end if
455 max_iter = this%max_iter
456
457 associate( w_d => this%w_d, r_d => this%r_d, d_d => this%d_d, &
458 blst => amg%blst)
459
460 if (.not. zero_initial_guess) then
461 call amg%device_matvec(this%w, x, w_d, x_d, this%lvl)
462 end if
463
464 thet = this%tha
465 delt = this%dlt
466 s1 = thet / delt
467 rhok = 1.0_rp / s1
468
469 ! First iteration
470 tmp1 = 1.0_rp / thet
471 call amg_device_cheby_solve_part1(r_d, f_d, w_d, x_d, d_d, &
472 tmp1, n, zero_initial_guess)
473 ! Rest of iterations
474 do iter = 2, max_iter
475 call amg%device_matvec(this%w, this%d, w_d, d_d, this%lvl)
476
477 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
478 tmp1 = rhokp1 * rhok
479 tmp2 = 2.0_rp * rhokp1 / delt
480 rhok = rhokp1
481
482 call amg_device_cheby_solve_part2(r_d, w_d, d_d, x_d, tmp1, tmp2, n)
483
484 end do
485 end associate
486 end subroutine amg_device_cheby_solve
487
493 subroutine amg_jacobi_init(this, n, lvl, max_iter)
494 class(amg_jacobi_t), intent(inout), target :: this
495 integer, intent(in) :: n
496 integer, intent(in) :: lvl
497 integer, intent(in) :: max_iter
498
499 allocate(this%d(n))
500 allocate(this%w(n))
501 allocate(this%r(n))
502 this%n = n
503 this%lvl = lvl
504 this%max_iter = max_iter
505 this%omega = 0.7_rp
506
507 end subroutine amg_jacobi_init
508
510 subroutine amg_jacobi_free(this)
511 class(amg_jacobi_t), intent(inout), target :: this
512 if (allocated(this%d)) then
513 deallocate(this%d)
514 end if
515 if (allocated(this%w)) then
516 deallocate(this%w)
517 end if
518 if (allocated(this%r)) then
519 deallocate(this%r)
520 end if
521 end subroutine amg_jacobi_free
522
526 subroutine amg_jacobi_diag(this, amg, n)
527 class(amg_jacobi_t), intent(inout) :: this
528 type(tamg_hierarchy_t), intent(inout) :: amg
529 integer, intent(in) :: n
530 real(kind=rp) :: val
531 integer :: i
532 do i = 1, n
533 call tamg_sample_matrix_val(val, amg, this%lvl, i, i)
534 this%d(i) = 1.0_rp / val
535 end do
536 this%recompute_diag = .false.
537 end subroutine amg_jacobi_diag
538
544 subroutine amg_jacobi_solve(this, x, f, n, amg, niter)
545 class(amg_jacobi_t), intent(inout) :: this
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
554 integer :: i
555
556 if (this%recompute_diag) then
557 call this%comp_diag(amg, n)
558 end if
559
560 if (present(niter)) then
561 max_iter = niter
562 else
563 max_iter = this%max_iter
564 end if
565
566 ! x = x + omega * Dinv( f - Ax )
567 associate( w => this%w, r => this%r, d => this%d)
568 do iter = 1, max_iter
569 w = 0.0_rp
571 call amg%matvec(w, x, this%lvl)
572 !$omp parallel private(i)
574 !$omp do
575 do i = 1, n
576 r(i) = f(i) - w(i)
577 end do
578 !$omp end do
580 !$omp do
581 do i = 1, n
582 r(i) = r(i) * d(i)
583 end do
584 !$omp end do
586 !$omp do
587 do i = 1, n
588 x(i) = x(i) + this%omega * r(i)
589 end do
590 !$omp end do
591 !$omp end parallel
592 end do
593 end associate
594 end subroutine amg_jacobi_solve
595
596 subroutine amg_smoo_monitor(lvl, smoo)
597 integer, intent(in) :: lvl
598 class(amg_cheby_t), intent(in) :: smoo
599 character(len=LOG_SIZE) :: log_buf
600
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)
605 end subroutine amg_smoo_monitor
606
607 subroutine amg_cheby_monitor(lvl, lam)
608 integer, intent(in) :: lvl
609 real(kind=rp), intent(in) :: lam
610 character(len=LOG_SIZE) :: log_buf
611
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)
615 end subroutine amg_cheby_monitor
616
617end module tree_amg_smoother
__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)
Definition device.F90:83
Copy data between host and device (or device and device)
Definition device.F90:72
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Defines a list of bc_t.
Definition bc_list.f90:34
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.
Definition device.F90:34
integer, parameter, public host_to_device
Definition device.F90:48
Gather-scatter.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
Logging routines.
Definition log.f90:34
type(log_t), public neko_log
Global log stream.
Definition log.f90:80
integer, parameter, public log_size
Definition log.f90:46
Definition math.f90:60
subroutine, public cmult(a, c, n)
Multiplication by constant c .
Definition math.f90:504
subroutine, public cmult2(a, b, c, n)
Multiplication by constant c .
Definition math.f90:519
real(kind=rp) function, public glsc3(a, b, c, n)
Weighted inner product .
Definition math.f90:1287
real(kind=rp) function, public glsc2(a, b, n)
Weighted inner product .
Definition math.f90:1266
subroutine, public add2(a, b, n)
Vector addition .
Definition math.f90:900
subroutine, public add3s2(a, b, c, c1, c2, n)
Returns .
Definition math.f90:1093
subroutine, public col2(a, b, n)
Vector multiplication .
Definition math.f90:1046
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:291
subroutine, public sub2(a, b, n)
Vector substraction .
Definition math.f90:948
subroutine, public add2s2(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on second argument)
Definition math.f90:998
Build configurations.
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Profiling interface.
Definition profiler.F90:34
subroutine, public profiler_start_region(name, region_id)
Started a named (name) profiler region.
Definition profiler.F90:79
subroutine, public profiler_end_region(name, region_id)
End the most recently started profiler region.
Definition profiler.F90:116
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.
Definition tree_amg.f90:34
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Gather-scatter kernel.
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:56
Type for a TreeAMG hierarchy.
Definition tree_amg.f90:100
Type for Chebyshev iteration using TreeAMG matvec.
Type for Chebyshev iteration using TreeAMG matvec.