Neko 1.99.5
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
51 use, intrinsic :: iso_c_binding
52 implicit none
53 private
54
56 type, public :: amg_jacobi_t
57 real(kind=rp), allocatable :: d(:)
58 real(kind=rp), allocatable :: w(:)
59 real(kind=rp), allocatable :: r(:)
60 real(kind=rp) :: omega
61 integer :: lvl
62 integer :: n
63 integer :: max_iter = 10
64 logical :: recompute_diag = .true.
65 contains
66 procedure, pass(this) :: init => amg_jacobi_init
67 procedure, pass(this) :: solve => amg_jacobi_solve
68 procedure, pass(this) :: comp_diag => amg_jacobi_diag
69 procedure, pass(this) :: free => amg_jacobi_free
70 end type amg_jacobi_t
71
73 type, public :: amg_cheby_t
74 real(kind=rp), allocatable :: d(:)
75 type(c_ptr) :: d_d = c_null_ptr
76 real(kind=rp), allocatable :: w(:)
77 type(c_ptr) :: w_d = c_null_ptr
78 real(kind=rp), allocatable :: r(:)
79 type(c_ptr) :: r_d = c_null_ptr
80 real(kind=rp) :: tha, dlt
81 integer :: lvl
82 integer :: n
83 integer :: power_its = 250
84 integer :: max_iter = 10
85 logical :: recompute_eigs = .true.
86 contains
87 procedure, pass(this) :: init => amg_cheby_init
88 procedure, pass(this) :: solve => amg_cheby_solve
89 procedure, pass(this) :: comp_eig => amg_cheby_power
90 procedure, pass(this) :: device_solve => amg_device_cheby_solve
91 procedure, pass(this) :: device_comp_eig => amg_device_cheby_power
92 procedure, pass(this) :: free => amg_cheby_free
93 end type amg_cheby_t
94
95contains
96
102 subroutine amg_cheby_init(this, n, lvl, max_iter)
103 class(amg_cheby_t), intent(inout), target :: this
104 integer, intent(in) :: n
105 integer, intent(in) :: lvl
106 integer, intent(in) :: max_iter
107
108 allocate(this%d(n))
109 allocate(this%w(n))
110 allocate(this%r(n))
111 if (neko_bcknd_device .eq. 1) then
112 call device_map(this%d, this%d_d, n)
113 call device_map(this%w, this%w_d, n)
114 call device_map(this%r, this%r_d, n)
115 end if
116 this%n = n
117 this%lvl = lvl
118 this%max_iter = max_iter
119 this%recompute_eigs = .true.
120
121 call amg_smoo_monitor(lvl, this)
122
123 end subroutine amg_cheby_init
124
126 subroutine amg_cheby_free(this)
127 class(amg_cheby_t), intent(inout), target :: this
128 if (allocated(this%d)) then
129 if (neko_bcknd_device .eq. 1 .and. c_associated(this%d_d)) then
130 call device_unmap(this%d, this%d_d)
131 end if
132 deallocate(this%d)
133 end if
134 if (allocated(this%w)) then
135 if (neko_bcknd_device .eq. 1 .and. c_associated(this%w_d)) then
136 call device_unmap(this%w, this%w_d)
137 end if
138 deallocate(this%w)
139 end if
140 if (allocated(this%r)) then
141 if (neko_bcknd_device .eq. 1 .and. c_associated(this%r_d)) then
142 call device_unmap(this%r, this%r_d)
143 end if
144 deallocate(this%r)
145 end if
146 end subroutine amg_cheby_free
147
148
152 subroutine amg_cheby_power(this, amg, n)
153 class(amg_cheby_t), intent(inout) :: this
154 type(tamg_hierarchy_t), intent(inout) :: amg
155 integer, intent(in) :: n
156 real(kind=rp) :: lam, b, a, rn
157 real(kind=rp), parameter :: boost = 1.1_rp
158 real(kind=rp), parameter :: lam_factor = 30.0_rp
159 real(kind=rp) :: wtw, dtw, dtd
160 integer, allocatable :: fixed_seed(:), saved_seed(:)
161 integer :: i, rnd_n
162 associate(w => this%w, d => this%d, coef => amg%coef, gs_h => amg%gs_h, &
163 msh => amg%msh, xh => amg%Xh, blst => amg%blst)
164
165 ! Save current random seed and set a fixed seed
166 call random_seed( size=rnd_n )
167 allocate(saved_seed(rnd_n))
168 allocate(fixed_seed(rnd_n))
169 fixed_seed = 3901
170 call random_seed( get=saved_seed )
171 call random_seed( put=fixed_seed )
172
173 do i = 1, n
174 call random_number(rn)
175 d(i) = rn + 10.0_rp
176 end do
177
178 ! Restore saved random seed
179 call random_seed( put=saved_seed )
180
181 if (this%lvl .eq. 0) then
182 call gs_h%op(d, n, gs_op_add)!TODO
183 call blst%apply(d, n)
184 end if
185 !Power method to get lamba max
186 do i = 1, this%power_its
187 call amg%matvec(w, d, this%lvl)
188
189 if (this%lvl .eq. 0) then
190 wtw = glsc3(w, coef%mult, w, n)
191 else
192 wtw = glsc2(w, w, n)
193 end if
194
195 call cmult2(d, w, 1.0_rp/sqrt(wtw), n)
196 end do
197
198 call amg%matvec(w, d, this%lvl)
199
200 if (this%lvl .eq. 0) then
201 dtw = glsc3(d, coef%mult, w, n)
202 dtd = glsc3(d, coef%mult, d, n)
203 else
204 dtw = glsc2(d, w, n)
205 dtd = glsc2(d, d, n)
206 end if
207 lam = dtw / dtd
208 b = lam * boost
209 a = lam / lam_factor
210 this%tha = (b+a)/2.0_rp
211 this%dlt = (b-a)/2.0_rp
212
213 this%recompute_eigs = .false.
214 call amg_cheby_monitor(this%lvl, lam)
215 end associate
216 end subroutine amg_cheby_power
217
224 subroutine amg_cheby_solve(this, x, f, n, amg, zero_init)
225 class(amg_cheby_t), intent(inout) :: this
226 integer, intent(in) :: n
227 real(kind=rp), dimension(n), intent(inout) :: x
228 real(kind=rp), dimension(n), intent(inout) :: f
229 class(tamg_hierarchy_t), intent(inout) :: amg
230 type(ksp_monitor_t) :: ksp_results
231 logical, optional, intent(in) :: zero_init
232 integer :: iter, max_iter, i
233 real(kind=rp) :: rtr, rnorm
234 real(kind=rp) :: rhok, rhokp1, s1, thet, delt, tmp1, tmp2
235 logical :: zero_initial_guess
236
237 if (this%recompute_eigs) then
238 call this%comp_eig(amg, n)
239 end if
240 if (present(zero_init)) then
241 zero_initial_guess = zero_init
242 else
243 zero_initial_guess = .false.
244 end if
245 max_iter = this%max_iter
246
247 associate( w => this%w, r => this%r, d => this%d, blst => amg%blst)
248 call copy(r, f, n)
249 if (.not. zero_initial_guess) then
250 call amg%matvec(w, x, this%lvl)
251 call sub2(r, w, n)
252 end if
253
254 thet = this%tha
255 delt = this%dlt
256 s1 = thet / delt
257 rhok = 1.0_rp / s1
258
259 ! First iteration
260 !OCL NORECURRENCE, NOVREC, NOALIAS
261 !DIR$ CONCURRENT
262 !GCC$ ivdep
263 !$omp parallel do
264 do i = 1, n
265 d(i) = 1.0_rp/thet * r(i)
266 x(i) = x(i) + d(i)
267 end do
268 !$omp end parallel do
269
270 ! Rest of iterations
271 do iter = 2, max_iter
272 call amg%matvec(w, d, this%lvl)
273
274 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
275 tmp1 = rhokp1 * rhok
276 tmp2 = 2.0_rp * rhokp1 / delt
277 rhok = rhokp1
278
279 !$omp parallel private(i)
280 !OCL NORECURRENCE, NOVREC, NOALIAS
281 !DIR$ CONCURRENT
282 !GCC$ ivdep
283 !$omp do
284 do i = 1, n
285 r(i) = r(i) - w(i)
286 d(i) = tmp1 * d(i) + tmp2 * r(i)
287 x(i) = x(i) + d(i)
288 end do
289 !$omp end do
290 !$omp end parallel
291 end do
292 end associate
293 end subroutine amg_cheby_solve
294
298 subroutine amg_device_cheby_power(this, amg, n)
299 class(amg_cheby_t), intent(inout) :: this
300 type(tamg_hierarchy_t), intent(inout) :: amg
301 integer, intent(in) :: n
302 real(kind=rp) :: lam, b, a, rn
303 real(kind=rp), parameter :: boost = 1.1_rp
304 real(kind=rp), parameter :: lam_factor = 30.0_rp
305 real(kind=rp) :: wtw, dtw, dtd
306 integer, allocatable :: fixed_seed(:), saved_seed(:)
307 integer :: i, rnd_n
308 associate(w => this%w, d => this%d, coef => amg%coef, gs_h => amg%gs_h, &
309 msh => amg%msh, xh => amg%Xh, blst => amg%blst)
310
311 ! Save current random seed and set a fixed seed
312 call random_seed( size=rnd_n )
313 allocate(saved_seed(rnd_n))
314 allocate(fixed_seed(rnd_n))
315 fixed_seed = 3901
316 call random_seed( get=saved_seed )
317 call random_seed( put=fixed_seed )
318
319 do i = 1, n
320 call random_number(rn)
321 d(i) = rn + 10.0_rp
322 end do
323 call device_memcpy(this%d, this%d_d, n, host_to_device, .true.)
324
325 ! Restore saved random seed
326 call random_seed( put=saved_seed )
327
328 if (this%lvl .eq. 0) then
329 call gs_h%op(d, n, gs_op_add)!TODO
330 call blst%apply(d, n)
331 end if
332 do i = 1, this%power_its
333 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
334
335 if (this%lvl .eq. 0) then
336 wtw = device_glsc3(this%w_d, coef%mult_d, this%w_d, n)
337 else
338 wtw = device_glsc2(this%w_d, this%w_d, n)
339 end if
340
341 call device_cmult2(this%d_d, this%w_d, 1.0_rp/sqrt(wtw), n)
342 end do
343
344 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
345
346 if (this%lvl .eq. 0) then
347 dtw = device_glsc3(this%d_d, coef%mult_d, this%w_d, n)
348 dtd = device_glsc3(this%d_d, coef%mult_d, this%d_d, n)
349 else
350 dtw = device_glsc2(this%d_d, this%w_d, n)
351 dtd = device_glsc2(this%d_d, this%d_d, n)
352 end if
353 lam = dtw / dtd
354 b = lam * boost
355 a = lam / lam_factor
356 this%tha = (b+a)/2.0_rp
357 this%dlt = (b-a)/2.0_rp
358
359 this%recompute_eigs = .false.
360 call amg_cheby_monitor(this%lvl, lam)
361 end associate
362 end subroutine amg_device_cheby_power
363
370 subroutine amg_device_cheby_solve(this, x, f, x_d, f_d, n, amg, zero_init)
371 class(amg_cheby_t), intent(inout) :: this
372 integer, intent(in) :: n
373 real(kind=rp), dimension(n), intent(inout) :: x
374 real(kind=rp), dimension(n), intent(inout) :: f
375 type(c_ptr) :: x_d
376 type(c_ptr) :: f_d
377 class(tamg_hierarchy_t), intent(inout) :: amg
378 type(ksp_monitor_t) :: ksp_results
379 logical, optional, intent(in) :: zero_init
380 integer :: iter, max_iter
381 real(kind=rp) :: rtr, rnorm
382 real(kind=rp) :: rhok, rhokp1, s1, thet, delt, tmp1, tmp2
383 logical :: zero_initial_guess
384
385 if (this%recompute_eigs) then
386 call this%device_comp_eig(amg, n)
387 end if
388 if (present(zero_init)) then
389 zero_initial_guess = zero_init
390 else
391 zero_initial_guess = .false.
392 end if
393 max_iter = this%max_iter
394
395 associate( w_d => this%w_d, r_d => this%r_d, d_d => this%d_d, &
396 blst => amg%blst)
397
398 if (.not. zero_initial_guess) then
399 call amg%device_matvec(this%w, x, w_d, x_d, this%lvl)
400 end if
401
402 thet = this%tha
403 delt = this%dlt
404 s1 = thet / delt
405 rhok = 1.0_rp / s1
406
407 ! First iteration
408 tmp1 = 1.0_rp / thet
409 call amg_device_cheby_solve_part1(r_d, f_d, w_d, x_d, d_d, &
410 tmp1, n, zero_initial_guess)
411 ! Rest of iterations
412 do iter = 2, max_iter
413 call amg%device_matvec(this%w, this%d, w_d, d_d, this%lvl)
414
415 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
416 tmp1 = rhokp1 * rhok
417 tmp2 = 2.0_rp * rhokp1 / delt
418 rhok = rhokp1
419
420 call amg_device_cheby_solve_part2(r_d, w_d, d_d, x_d, tmp1, tmp2, n)
421
422 end do
423 end associate
424 end subroutine amg_device_cheby_solve
425
431 subroutine amg_jacobi_init(this, n, lvl, max_iter)
432 class(amg_jacobi_t), intent(inout), target :: this
433 integer, intent(in) :: n
434 integer, intent(in) :: lvl
435 integer, intent(in) :: max_iter
436
437 allocate(this%d(n))
438 allocate(this%w(n))
439 allocate(this%r(n))
440 this%n = n
441 this%lvl = lvl
442 this%max_iter = max_iter
443 this%omega = 0.7_rp
444
445 end subroutine amg_jacobi_init
446
448 subroutine amg_jacobi_free(this)
449 class(amg_jacobi_t), intent(inout), target :: this
450 if (allocated(this%d)) then
451 deallocate(this%d)
452 end if
453 if (allocated(this%w)) then
454 deallocate(this%w)
455 end if
456 if (allocated(this%r)) then
457 deallocate(this%r)
458 end if
459 end subroutine amg_jacobi_free
460
464 subroutine amg_jacobi_diag(this, amg, n)
465 class(amg_jacobi_t), intent(inout) :: this
466 type(tamg_hierarchy_t), intent(inout) :: amg
467 integer, intent(in) :: n
468 real(kind=rp) :: val
469 integer :: i
470 do i = 1, n
471 call tamg_sample_matrix_val(val, amg, this%lvl, i, i)
472 this%d(i) = 1.0_rp / val
473 end do
474 this%recompute_diag = .false.
475 end subroutine amg_jacobi_diag
476
482 subroutine amg_jacobi_solve(this, x, f, n, amg, niter)
483 class(amg_jacobi_t), intent(inout) :: this
484 integer, intent(in) :: n
485 real(kind=rp), dimension(n), intent(inout) :: x
486 real(kind=rp), dimension(n), intent(inout) :: f
487 class(tamg_hierarchy_t), intent(inout) :: amg
488 type(ksp_monitor_t) :: ksp_results
489 integer, optional, intent(in) :: niter
490 integer :: iter, max_iter
491 real(kind=rp) :: rtr, rnorm
492 integer :: i
493
494 if (this%recompute_diag) then
495 call this%comp_diag(amg, n)
496 end if
497
498 if (present(niter)) then
499 max_iter = niter
500 else
501 max_iter = this%max_iter
502 end if
503
504 ! x = x + omega * Dinv( f - Ax )
505 associate( w => this%w, r => this%r, d => this%d)
506 do iter = 1, max_iter
507 w = 0.0_rp
509 call amg%matvec(w, x, this%lvl)
510 !$omp parallel private(i)
512 !$omp do
513 do i = 1, n
514 r(i) = f(i) - w(i)
515 end do
516 !$omp end do
518 !$omp do
519 do i = 1, n
520 r(i) = r(i) * d(i)
521 end do
522 !$omp end do
524 !$omp do
525 do i = 1, n
526 x(i) = x(i) + this%omega * r(i)
527 end do
528 !$omp end do
529 !$omp end parallel
530 end do
531 end associate
532 end subroutine amg_jacobi_solve
533
534 subroutine amg_smoo_monitor(lvl, smoo)
535 integer, intent(in) :: lvl
536 class(amg_cheby_t), intent(in) :: smoo
537 character(len=LOG_SIZE) :: log_buf
538
539 write(log_buf, '(A8,I2,A28)') '-- level', lvl, '-- init smoother: Chebyshev'
540 call neko_log%message(log_buf)
541 write(log_buf, '(A22,I6)') 'Iterations:', smoo%max_iter
542 call neko_log%message(log_buf)
543 end subroutine amg_smoo_monitor
544
545 subroutine amg_cheby_monitor(lvl, lam)
546 integer, intent(in) :: lvl
547 real(kind=rp), intent(in) :: lam
548 character(len=LOG_SIZE) :: log_buf
549
550 write(log_buf, '(A12,I2,A29,F12.3)') '-- AMG level', lvl, &
551 '-- Chebyshev approx. max eig', lam
552 call neko_log%message(log_buf)
553 end subroutine amg_cheby_monitor
554
555end 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:78
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:84
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
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
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:56
Type for a TreeAMG hierarchy.
Definition tree_amg.f90:87
Type for Chebyshev iteration using TreeAMG matvec.
Type for Chebyshev iteration using TreeAMG matvec.