Neko 1.99.6
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 !DIR$ IVDEP
263 !GCC$ ivdep
264 !$omp parallel do
265 do i = 1, n
266 d(i) = 1.0_rp/thet * r(i)
267 x(i) = x(i) + d(i)
268 end do
269 !$omp end parallel do
270
271 ! Rest of iterations
272 do iter = 2, max_iter
273 call amg%matvec(w, d, this%lvl)
274
275 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
276 tmp1 = rhokp1 * rhok
277 tmp2 = 2.0_rp * rhokp1 / delt
278 rhok = rhokp1
279
280 !$omp parallel private(i)
281 !OCL NORECURRENCE, NOVREC, NOALIAS
282 !DIR$ CONCURRENT
283 !DIR$ IVDEP
284 !GCC$ ivdep
285 !$omp do
286 do i = 1, n
287 r(i) = r(i) - w(i)
288 d(i) = tmp1 * d(i) + tmp2 * r(i)
289 x(i) = x(i) + d(i)
290 end do
291 !$omp end do
292 !$omp end parallel
293 end do
294 end associate
295 end subroutine amg_cheby_solve
296
300 subroutine amg_device_cheby_power(this, amg, n)
301 class(amg_cheby_t), intent(inout) :: this
302 type(tamg_hierarchy_t), intent(inout) :: amg
303 integer, intent(in) :: n
304 real(kind=rp) :: lam, b, a, rn
305 real(kind=rp), parameter :: boost = 1.1_rp
306 real(kind=rp), parameter :: lam_factor = 30.0_rp
307 real(kind=rp) :: wtw, dtw, dtd
308 integer, allocatable :: fixed_seed(:), saved_seed(:)
309 integer :: i, rnd_n
310 associate(w => this%w, d => this%d, coef => amg%coef, gs_h => amg%gs_h, &
311 msh => amg%msh, xh => amg%Xh, blst => amg%blst)
312
313 ! Save current random seed and set a fixed seed
314 call random_seed( size=rnd_n )
315 allocate(saved_seed(rnd_n))
316 allocate(fixed_seed(rnd_n))
317 fixed_seed = 3901
318 call random_seed( get=saved_seed )
319 call random_seed( put=fixed_seed )
320
321 do i = 1, n
322 call random_number(rn)
323 d(i) = rn + 10.0_rp
324 end do
325 call device_memcpy(this%d, this%d_d, n, host_to_device, .true.)
326
327 ! Restore saved random seed
328 call random_seed( put=saved_seed )
329
330 if (this%lvl .eq. 0) then
331 call gs_h%op(d, n, gs_op_add)!TODO
332 call blst%apply(d, n)
333 end if
334 do i = 1, this%power_its
335 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
336
337 if (this%lvl .eq. 0) then
338 wtw = device_glsc3(this%w_d, coef%mult_d, this%w_d, n)
339 else
340 wtw = device_glsc2(this%w_d, this%w_d, n)
341 end if
342
343 call device_cmult2(this%d_d, this%w_d, 1.0_rp/sqrt(wtw), n)
344 end do
345
346 call amg%device_matvec(w, d, this%w_d, this%d_d, this%lvl)
347
348 if (this%lvl .eq. 0) then
349 dtw = device_glsc3(this%d_d, coef%mult_d, this%w_d, n)
350 dtd = device_glsc3(this%d_d, coef%mult_d, this%d_d, n)
351 else
352 dtw = device_glsc2(this%d_d, this%w_d, n)
353 dtd = device_glsc2(this%d_d, this%d_d, n)
354 end if
355 lam = dtw / dtd
356 b = lam * boost
357 a = lam / lam_factor
358 this%tha = (b+a)/2.0_rp
359 this%dlt = (b-a)/2.0_rp
360
361 this%recompute_eigs = .false.
362 call amg_cheby_monitor(this%lvl, lam)
363 end associate
364 end subroutine amg_device_cheby_power
365
372 subroutine amg_device_cheby_solve(this, x, f, x_d, f_d, n, amg, zero_init)
373 class(amg_cheby_t), intent(inout) :: this
374 integer, intent(in) :: n
375 real(kind=rp), dimension(n), intent(inout) :: x
376 real(kind=rp), dimension(n), intent(inout) :: f
377 type(c_ptr) :: x_d
378 type(c_ptr) :: f_d
379 class(tamg_hierarchy_t), intent(inout) :: amg
380 type(ksp_monitor_t) :: ksp_results
381 logical, optional, intent(in) :: zero_init
382 integer :: iter, max_iter
383 real(kind=rp) :: rtr, rnorm
384 real(kind=rp) :: rhok, rhokp1, s1, thet, delt, tmp1, tmp2
385 logical :: zero_initial_guess
386
387 if (this%recompute_eigs) then
388 call this%device_comp_eig(amg, n)
389 end if
390 if (present(zero_init)) then
391 zero_initial_guess = zero_init
392 else
393 zero_initial_guess = .false.
394 end if
395 max_iter = this%max_iter
396
397 associate( w_d => this%w_d, r_d => this%r_d, d_d => this%d_d, &
398 blst => amg%blst)
399
400 if (.not. zero_initial_guess) then
401 call amg%device_matvec(this%w, x, w_d, x_d, this%lvl)
402 end if
403
404 thet = this%tha
405 delt = this%dlt
406 s1 = thet / delt
407 rhok = 1.0_rp / s1
408
409 ! First iteration
410 tmp1 = 1.0_rp / thet
411 call amg_device_cheby_solve_part1(r_d, f_d, w_d, x_d, d_d, &
412 tmp1, n, zero_initial_guess)
413 ! Rest of iterations
414 do iter = 2, max_iter
415 call amg%device_matvec(this%w, this%d, w_d, d_d, this%lvl)
416
417 rhokp1 = 1.0_rp / (2.0_rp * s1 - rhok)
418 tmp1 = rhokp1 * rhok
419 tmp2 = 2.0_rp * rhokp1 / delt
420 rhok = rhokp1
421
422 call amg_device_cheby_solve_part2(r_d, w_d, d_d, x_d, tmp1, tmp2, n)
423
424 end do
425 end associate
426 end subroutine amg_device_cheby_solve
427
433 subroutine amg_jacobi_init(this, n, lvl, max_iter)
434 class(amg_jacobi_t), intent(inout), target :: this
435 integer, intent(in) :: n
436 integer, intent(in) :: lvl
437 integer, intent(in) :: max_iter
438
439 allocate(this%d(n))
440 allocate(this%w(n))
441 allocate(this%r(n))
442 this%n = n
443 this%lvl = lvl
444 this%max_iter = max_iter
445 this%omega = 0.7_rp
446
447 end subroutine amg_jacobi_init
448
450 subroutine amg_jacobi_free(this)
451 class(amg_jacobi_t), intent(inout), target :: this
452 if (allocated(this%d)) then
453 deallocate(this%d)
454 end if
455 if (allocated(this%w)) then
456 deallocate(this%w)
457 end if
458 if (allocated(this%r)) then
459 deallocate(this%r)
460 end if
461 end subroutine amg_jacobi_free
462
466 subroutine amg_jacobi_diag(this, amg, n)
467 class(amg_jacobi_t), intent(inout) :: this
468 type(tamg_hierarchy_t), intent(inout) :: amg
469 integer, intent(in) :: n
470 real(kind=rp) :: val
471 integer :: i
472 do i = 1, n
473 call tamg_sample_matrix_val(val, amg, this%lvl, i, i)
474 this%d(i) = 1.0_rp / val
475 end do
476 this%recompute_diag = .false.
477 end subroutine amg_jacobi_diag
478
484 subroutine amg_jacobi_solve(this, x, f, n, amg, niter)
485 class(amg_jacobi_t), intent(inout) :: this
486 integer, intent(in) :: n
487 real(kind=rp), dimension(n), intent(inout) :: x
488 real(kind=rp), dimension(n), intent(inout) :: f
489 class(tamg_hierarchy_t), intent(inout) :: amg
490 type(ksp_monitor_t) :: ksp_results
491 integer, optional, intent(in) :: niter
492 integer :: iter, max_iter
493 real(kind=rp) :: rtr, rnorm
494 integer :: i
495
496 if (this%recompute_diag) then
497 call this%comp_diag(amg, n)
498 end if
499
500 if (present(niter)) then
501 max_iter = niter
502 else
503 max_iter = this%max_iter
504 end if
505
506 ! x = x + omega * Dinv( f - Ax )
507 associate( w => this%w, r => this%r, d => this%d)
508 do iter = 1, max_iter
509 w = 0.0_rp
511 call amg%matvec(w, x, this%lvl)
512 !$omp parallel private(i)
514 !$omp do
515 do i = 1, n
516 r(i) = f(i) - w(i)
517 end do
518 !$omp end do
520 !$omp do
521 do i = 1, n
522 r(i) = r(i) * d(i)
523 end do
524 !$omp end do
526 !$omp do
527 do i = 1, n
528 x(i) = x(i) + this%omega * r(i)
529 end do
530 !$omp end do
531 !$omp end parallel
532 end do
533 end associate
534 end subroutine amg_jacobi_solve
535
536 subroutine amg_smoo_monitor(lvl, smoo)
537 integer, intent(in) :: lvl
538 class(amg_cheby_t), intent(in) :: smoo
539 character(len=LOG_SIZE) :: log_buf
540
541 write(log_buf, '(A8,I2,A28)') '-- level', lvl, '-- init smoother: Chebyshev'
542 call neko_log%message(log_buf)
543 write(log_buf, '(A22,I6)') 'Iterations:', smoo%max_iter
544 call neko_log%message(log_buf)
545 end subroutine amg_smoo_monitor
546
547 subroutine amg_cheby_monitor(lvl, lam)
548 integer, intent(in) :: lvl
549 real(kind=rp), intent(in) :: lam
550 character(len=LOG_SIZE) :: log_buf
551
552 write(log_buf, '(A12,I2,A29,F12.3)') '-- AMG level', lvl, &
553 '-- Chebyshev approx. max eig', lam
554 call neko_log%message(log_buf)
555 end subroutine amg_cheby_monitor
556
557end 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
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:87
Type for Chebyshev iteration using TreeAMG matvec.
Type for Chebyshev iteration using TreeAMG matvec.