Neko 1.99.7
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
tree_amg_multigrid.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!
45 use num_types, only: rp
46 use utils, only : neko_error, neko_warning
47 use math, only : add2, rzero, glsc2, col2, copy, add2s1
50 use comm
51 use mpi_f08, only: mpi_allreduce, mpi_min, mpi_in_place, mpi_integer
52 use coefs, only : coef_t
53 use mesh, only : mesh_t
54 use space, only : space_t
55 use ax_product, only: ax_t
56 use bc_list, only : bc_list_t
57 use gather_scatter, only : gs_t, gs_op_add
63 use logger, only : neko_log, log_size
67 use, intrinsic :: iso_c_binding
68 !$ use omp_lib, only : omp_get_max_threads
69 implicit none
70 private
71
72 type :: tamg_wrk_t
73 integer :: n = -1
74 real(kind=rp), allocatable :: r(:)
75 real(kind=rp), allocatable :: b(:)
76 real(kind=rp), allocatable :: x(:)
77 type(c_ptr) :: r_d = c_null_ptr
78 type(c_ptr) :: b_d = c_null_ptr
79 type(c_ptr) :: x_d = c_null_ptr
80 end type tamg_wrk_t
81
83 type, public :: tamg_solver_t
84 type(tamg_hierarchy_t), allocatable :: amg
85 type(amg_cheby_t), allocatable :: smoo(:)
86 type(tamg_wrk_t), allocatable :: wrk(:)
87 !type(amg_jacobi_t), allocatable :: jsmoo(:)
88 integer :: nlvls
89 integer :: max_iter
90 contains
91 procedure, pass(this) :: init => tamg_mg_init
92 procedure, pass(this) :: solve => tamg_mg_solve
93 procedure, pass(this) :: free => tamg_mg_free
94 procedure, pass(this) :: invalidate_eigs => tamg_mg_invalidate_eigs
95 procedure, pass(this) :: set_eig_refresh => tamg_mg_set_eig_refresh
96 procedure, private, pass(this) :: mg_cycle => tamg_mg_cycle
97 procedure, private, pass(this) :: mg_cycle_d => tamg_mg_cycle_d
98 end type tamg_solver_t
99
100contains
101
111 subroutine tamg_mg_init(this, ax, Xh, coef, msh, gs_h, nlvls, blst, &
112 max_iter, cheby_degree)
113 class(tamg_solver_t), intent(inout), target :: this
114 class(ax_t), target, intent(in) :: ax
115 type(space_t), target, intent(in) :: Xh
116 type(coef_t), target, intent(in) :: coef
117 type(mesh_t), target, intent(in) :: msh
118 type(gs_t), target, intent(in) :: gs_h
119 type(bc_list_t), target, intent(in) :: blst
120 integer, intent(in) :: nlvls
121 integer, intent(in) :: max_iter
122 integer, intent(in) :: cheby_degree
123 integer :: lvl, n, mlvl, target_num_aggs
124 integer, allocatable :: agg_nhbr(:,:), nhbr_tmp(:,:)
125 character(len=LOG_SIZE) :: log_buf
126 integer :: glb_min_target_aggs
127 logical :: use_greedy_agg
128
129 call neko_log%section('AMG')
130
131 write(log_buf, '(A28,I2,A8)') 'Creating AMG hierarchy with', &
132 nlvls, 'levels.'
133 call neko_log%message(log_buf)
134
135 allocate( this%amg )
136 call this%amg%init(ax, xh, coef, msh, gs_h, nlvls, blst)
137
138 ! Aggregation
139 use_greedy_agg = .true.
140 ! Create level 1 (neko elements are level 0)
141 call aggregate_finest_level(this%amg, xh%lx, xh%ly, xh%lz, msh%nelv)
142
143 ! Create the remaining levels
144 allocate( agg_nhbr, source = msh%facet_neigh )
145 do mlvl = 2, nlvls-1
146 ! estimate number of aggregates
147 if (use_greedy_agg) then
148 target_num_aggs = this%amg%lvl(mlvl-1)%nnodes / 8
149 else
150 target_num_aggs = this%amg%lvl(mlvl-1)%nnodes / 2
151 end if
152
153 glb_min_target_aggs = target_num_aggs
154 call mpi_allreduce(mpi_in_place, glb_min_target_aggs, 1, &
155 mpi_integer, mpi_min, neko_comm)
156 if (glb_min_target_aggs .lt. 4 ) then
157 call neko_warning( &
158 "TAMG: Too many levels. Not enough DOFs for coarsest grid.")
159 this%amg%nlvls = mlvl
160 exit
161 end if
162
163 if (use_greedy_agg) then
164 call print_preagg_info( mlvl, glb_min_target_aggs, 1)
165 call aggregate_greedy(this%amg, mlvl, target_num_aggs, &
166 agg_nhbr, nhbr_tmp)
167 else
168 call print_preagg_info( mlvl, glb_min_target_aggs, 2)
169 call aggregate_pairs(this%amg, mlvl, target_num_aggs, &
170 agg_nhbr, nhbr_tmp)
171 end if
172
173 agg_nhbr = nhbr_tmp
174 deallocate( nhbr_tmp )
175 end do
176 deallocate( agg_nhbr )
177
178 ! Create the end point
179 call aggregate_end(this%amg, this%amg%nlvls)
180
181 this%max_iter = max_iter
182
183 this%nlvls = this%amg%nlvls
184 if (this%nlvls .gt. this%amg%nlvls) then
185 call neko_error( &
186 "Requested number multigrid levels &
187 & is greater than the initialized AMG levels")
188 end if
189
190 ! Initialize relaxation methods
191 allocate(this%smoo(0:(this%amg%nlvls)))
192 do lvl = 0, this%amg%nlvls-1
193 n = this%amg%lvl(lvl+1)%fine_lvl_dofs
194 call this%smoo(lvl)%init(n, lvl, cheby_degree)
195 end do
196
197 ! Allocate work space on each level
198 allocate(this%wrk(0:(this%amg%nlvls)))
199 do lvl = 0, this%amg%nlvls-1
200 n = this%amg%lvl(lvl+1)%fine_lvl_dofs
201 this%wrk(lvl)%n = n
202 allocate( this%wrk(lvl)%r(n) )
203 allocate( this%wrk(lvl)%b(n) )
204 allocate( this%wrk(lvl)%x(n) )
205 if (neko_bcknd_device .eq. 1) then
206 call device_map( this%wrk(lvl)%r, this%wrk(lvl)%r_d, n)
207 call device_map( this%wrk(lvl)%b, this%wrk(lvl)%b_d, n)
208 call device_map( this%wrk(lvl)%x, this%wrk(lvl)%x_d, n)
209 end if
210 end do
211
212 !allocate(this%jsmoo(0:(this%amg%nlvls)))
213 !do lvl = 0, this%amg%nlvls-1
214 ! n = this%amg%lvl(lvl+1)%fine_lvl_dofs
215 ! call this%jsmoo(lvl)%init(n ,lvl, cheby_degree)
216 !end do
217
218 ! Create index mapping between levels
219 call fill_lvl_map(this%amg)
220
221 ! Invert those maps so the flat matvec can be driven from the aggregate
222 ! side instead of scattering into the coarse vector with atomics
223 call build_agg_csr(this%amg)
224
225 call neko_log%end_section()
226
227 end subroutine tamg_mg_init
228
230 subroutine tamg_mg_free(this)
231 class(tamg_solver_t), intent(inout), target :: this
232 integer :: i
233 if (allocated(this%amg)) then
234 call this%amg%free()
235 deallocate(this%amg)
236 end if
237 if (allocated(this%smoo)) then
238 do i = 0, (size(this%smoo)-1)
239 call this%smoo(i)%free()
240 end do
241 deallocate(this%smoo)
242 end if
243 if (allocated(this%wrk)) then
244 do i = 0, (size(this%wrk)-1)
245 if (allocated(this%wrk(i)%r)) then
246 if (neko_bcknd_device .eq. 1 .and. &
247 c_associated(this%wrk(i)%r_d)) then
248 call device_unmap(this%wrk(i)%r, this%wrk(i)%r_d)
249 end if
250 deallocate(this%wrk(i)%r)
251 end if
252 if (allocated(this%wrk(i)%b)) then
253 if (neko_bcknd_device .eq. 1 .and. &
254 c_associated(this%wrk(i)%b_d)) then
255 call device_unmap(this%wrk(i)%b, this%wrk(i)%b_d)
256 end if
257 deallocate(this%wrk(i)%b)
258 end if
259 if (allocated(this%wrk(i)%x)) then
260 if (neko_bcknd_device .eq. 1 .and. &
261 c_associated(this%wrk(i)%x_d)) then
262 call device_unmap(this%wrk(i)%x, this%wrk(i)%x_d)
263 end if
264 deallocate(this%wrk(i)%x)
265 end if
266 end do
267 end if
268 end subroutine tamg_mg_free
269
270
273 subroutine tamg_mg_invalidate_eigs(this)
274 class(tamg_solver_t), intent(inout) :: this
275 integer :: lvl
276 do lvl = 0, this%amg%nlvls-1
277 this%smoo(lvl)%recompute_eigs = .true.
278 end do
279 end subroutine tamg_mg_invalidate_eigs
280
281
285 subroutine tamg_mg_set_eig_refresh(this, warm_start, power_its_refresh)
286 class(tamg_solver_t), intent(inout) :: this
287 logical, intent(in) :: warm_start
288 integer, intent(in) :: power_its_refresh
289 integer :: lvl
290 do lvl = 0, this%amg%nlvls-1
291 this%smoo(lvl)%warm_start_eigs = warm_start
292 this%smoo(lvl)%power_its_refresh = power_its_refresh
293 end do
294 end subroutine tamg_mg_set_eig_refresh
295
296
301 subroutine tamg_mg_solve(this, z, r, n)
302 integer, intent(in) :: n
303 class(tamg_solver_t), intent(inout) :: this
304 real(kind=rp), dimension(n), intent(inout) :: z
305 real(kind=rp), dimension(n), intent(inout) :: r
306 type(c_ptr) :: z_d
307 type(c_ptr) :: r_d
308 integer :: iter, max_iter, i
309 logical :: zero_initial_guess
310
311 max_iter = this%max_iter
312
313 if (neko_bcknd_device .eq. 1) then
314 z_d = device_get_ptr(z)
315 r_d = device_get_ptr(r)
316 ! Zero out the initial guess because we do not handle null
317 ! spaces very well...
318 call device_rzero(this%wrk(0)%x_d, n)
319 call device_copy(this%wrk(0)%b_d, r_d, n)
320 zero_initial_guess = .true.
321 ! Call the amg cycle
322 do iter = 1, max_iter
323 call this%mg_cycle_d(zero_initial_guess)
324 zero_initial_guess = .false.
325 end do
326 call device_copy(z_d, this%wrk(0)%x_d, n)
327 else
328 ! Zero out the initial guess becuase we do not handle null spaces
329 ! very well...
330 !OCL NORECURRENCE, NOVREC, NOALIAS
331 !DIR$ CONCURRENT
332 !DIR$ IVDEP
333 !GCC$ ivdep
334 !$omp parallel do
335 do i = 1, n
336 this%wrk(0)%x(i) = 0.0_rp
337 this%wrk(0)%b(i) = r(i)
338 end do
339 !$omp end parallel do
340 zero_initial_guess = .true.
341 ! Call the amg cycle
342 do iter = 1, max_iter
343 call this%mg_cycle(zero_initial_guess)
344 zero_initial_guess = .false.
345 end do
346 call copy(z, this%wrk(0)%x, n)
347 end if
348 end subroutine tamg_mg_solve
349
350
354 subroutine tamg_mg_cycle(this, zero_initial_guess)
355 class(tamg_solver_t), intent(inout), target :: this
356 logical, intent(inout) :: zero_initial_guess
357 character(len=2) :: lvl_name
358 integer :: max_lvl, lvl
359
360 max_lvl = this%nlvls-1
361 ! Loop down hierarchy. Fine to coarse
362 do lvl = 0, max_lvl-1
363 write(lvl_name, '(I0)') lvl
364 call profiler_start_region( "AMG_level_" // trim(lvl_name))
365 associate(x => this%wrk(lvl)%x, b => this%wrk(lvl)%b, &
366 r => this%wrk(lvl)%r, n => this%wrk(lvl)%n)
367 !!----------!!
368 !! SMOOTH !!
369 !!----------!!
370 call this%smoo(lvl)%solve(x, b, n, this%amg, &
371 zero_initial_guess)
372 !!----------!!
373 !! Residual !!
374 !!----------!!
375 call calc_resid(r, x, b, this%amg, lvl, n)
376 !!----------!!
377 !! Restrict !!
378 !!----------!!
379 call this%amg%interp_f2c(this%wrk(lvl+1)%b, r, lvl+1)
380
381 call rzero(this%wrk(lvl+1)%x, this%wrk(lvl+1)%n)
382 zero_initial_guess = .true.
383 end associate
384 call profiler_end_region( "AMG_level_" // trim(lvl_name))
385 end do
386 write(lvl_name, '(I0)') max_lvl
387 call profiler_start_region( "AMG_level_" // trim(lvl_name))
388 !!-------------------!!
389 !! Call Coarse solve !!
390 !!-------------------!!
391 call this%smoo(max_lvl)%solve(this%wrk(max_lvl)%x, &
392 this%wrk(max_lvl)%b, this%amg%lvl(max_lvl)%nnodes, this%amg, &
393 zero_initial_guess)
394 call profiler_end_region( "AMG_level_" // trim(lvl_name))
395
396 zero_initial_guess = .false.
397 ! Loop up hierarchy. Coarse to fine
398 do lvl = max_lvl-1, 0, -1
399 write(lvl_name, '(I0)') lvl
400 call profiler_start_region( "AMG_level_" // trim(lvl_name))
401 associate(x => this%wrk(lvl)%x, b => this%wrk(lvl)%b, &
402 r => this%wrk(lvl)%r, n => this%wrk(lvl)%n)
403 !!----------!!
404 !! Project !!
405 !!----------!!
406 call this%amg%interp_c2f(r, this%wrk(lvl+1)%x, lvl+1)
407 !!----------!!
408 !! Correct !!
409 !!----------!!
410 call add2(x, r, n)
411 !!----------!!
412 !! SMOOTH !!
413 !!----------!!
414 call this%smoo(lvl)%solve(x, b, n, this%amg)
415 end associate
416 call profiler_end_region( "AMG_level_" // trim(lvl_name))
417 end do
418 end subroutine tamg_mg_cycle
419
423 subroutine tamg_mg_cycle_d(this, zero_initial_guess)
424 class(tamg_solver_t), intent(inout), target :: this
425 logical, intent(inout) :: zero_initial_guess
426 character(len=2) :: lvl_name
427 integer :: max_lvl, lvl
428
429 max_lvl = this%nlvls-1
430 ! Loop down hierarchy. Fine to coarse
431 do lvl = 0, max_lvl-1
432 write(lvl_name, '(I0)') lvl
433 call profiler_start_region( "AMG_level_" // trim(lvl_name))
434 associate(x => this%wrk(lvl)%x, x_d => this%wrk(lvl)%x_d, &
435 b => this%wrk(lvl)%b, b_d => this%wrk(lvl)%b_d, &
436 r => this%wrk(lvl)%r, r_d => this%wrk(lvl)%r_d, &
437 n => this%wrk(lvl)%n)
438 !!----------!!
439 !! SMOOTH !!
440 !!----------!!
441 call this%smoo(lvl)%device_solve(x, b, x_d, b_d, n, this%amg, &
442 zero_initial_guess)
443 !!----------!!
444 !! Residual !!
445 !!----------!!
446 call this%amg%device_matvec(r, x, r_d, x_d, lvl)
447 call device_sub3(r_d, b_d, r_d, n)
448 !!----------!!
449 !! Restrict !!
450 !!----------!!
451 call this%amg%interp_f2c_d(this%wrk(lvl+1)%b_d, r_d, lvl+1)
452
453 call device_rzero(this%wrk(lvl+1)%x_d, this%wrk(lvl+1)%n)
454 zero_initial_guess = .true.
455 end associate
456 call profiler_end_region( "AMG_level_" // trim(lvl_name))
457 end do
458 write(lvl_name, '(I0)') max_lvl
459 call profiler_start_region( "AMG_level_" // trim(lvl_name))
460 !!-------------------!!
461 !! Call Coarse solve !!
462 !!-------------------!!
463 call this%smoo(max_lvl)%device_solve( &
464 this%wrk(max_lvl)%x, this%wrk(max_lvl)%b, &
465 this%wrk(max_lvl)%x_d, this%wrk(max_lvl)%b_d, &
466 this%amg%lvl(max_lvl)%nnodes, this%amg, &
467 zero_initial_guess)
468 call profiler_end_region( "AMG_level_" // trim(lvl_name))
469
470 zero_initial_guess = .false.
471 ! Loop up hierarchy. Coarse to fine
472 do lvl = max_lvl-1, 0, -1
473 write(lvl_name, '(I0)') lvl
474 call profiler_start_region( "AMG_level_" // trim(lvl_name))
475 associate(x => this%wrk(lvl)%x, x_d => this%wrk(lvl)%x_d, &
476 b => this%wrk(lvl)%b, b_d => this%wrk(lvl)%b_d, &
477 r => this%wrk(lvl)%r, r_d => this%wrk(lvl)%r_d, &
478 n => this%wrk(lvl)%n)
479 !!----------!!
480 !! Project !!
481 !!----------!!
482 call this%amg%interp_c2f_d(r_d, this%wrk(lvl+1)%x_d, lvl+1, r)
483 !!----------!!
484 !! Correct !!
485 !!----------!!
486 call device_add2(x_d, r_d, n)
487 !!----------!!
488 !! SMOOTH !!
489 !!----------!!
490 call this%smoo(lvl)%device_solve(x, b, x_d, b_d, n, this%amg)
491 end associate
492 call profiler_end_region( "AMG_level_" // trim(lvl_name))
493 end do
494 end subroutine tamg_mg_cycle_d
495
496
504 subroutine calc_resid(r, x, b, amg, lvl, n)
505 integer, intent(in) :: n
506 real(kind=rp), intent(inout) :: r(n)
507 real(kind=rp), intent(inout) :: x(n)
508 real(kind=rp), intent(inout) :: b(n)
509 type(tamg_hierarchy_t), intent(inout) :: amg
510 integer, intent(in) :: lvl
511 integer :: i
512 call amg%matvec(r, x, lvl)
513 call add2s1(r, b, -1.0_rp, n)
514 end subroutine calc_resid
515
516
517 subroutine print_preagg_info(lvl, nagg, agg_type)
518 integer, intent(in) :: lvl, nagg, agg_type
519 character(len=LOG_SIZE) :: log_buf
520 !TODO: calculate min and max agg size
521 if (agg_type .eq. 1) then
522 write(log_buf, '(A8,I2,A31)') '-- level', lvl, &
523 '-- Calling Greedy Aggregation'
524 else if (agg_type .eq. 2) then
525 write(log_buf, '(A8,I2,A33)') '-- level', lvl, &
526 '-- Calling Pairwise Aggregation'
527 else
528 write(log_buf, '(A8,I2,A31)') '-- level', lvl, &
529 '-- UNKNOWN Aggregation'
530 end if
531 call neko_log%message(log_buf)
532 write(log_buf, '(A33,I6)') 'Target Aggregates:', nagg
533 call neko_log%message(log_buf)
534 end subroutine print_preagg_info
535
536 subroutine print_resid_info(r, x, b, r_d, x_d, b_d, amg, lvl, n)
537 integer, intent(in) :: lvl, n
538 real(kind=rp), intent(inout) :: r(n)
539 real(kind=rp), intent(inout) :: x(n)
540 real(kind=rp), intent(inout) :: b(n)
541 type(c_ptr) :: r_d
542 type(c_ptr) :: x_d
543 type(c_ptr) :: b_d
544 type(tamg_hierarchy_t), intent(inout) :: amg
545 real(kind=rp) :: val
546 character(len=LOG_SIZE) :: log_buf
547
548 call amg%device_matvec(r, x, r_d, x_d, lvl)
549 call device_sub3(r_d, b_d, r_d, n)
550 val = device_glsc2(r_d, r_d, n)
551
552 write(log_buf, '(A33,I6,F12.6)') 'tAMG resid:', lvl, val
553 call neko_log%message(log_buf)
554 end subroutine print_resid_info
555
558 subroutine fill_lvl_map(amg)
559 type(tamg_hierarchy_t), intent(inout) :: amg
560 integer, allocatable :: dof2gid(:)
561 integer :: i, j, k, l, nid, n, nmap
562 do j = 1, amg%lvl(1)%nnodes
563 do k = 1, amg%lvl(1)%nodes(j)%ndofs
564 nid = amg%lvl(1)%nodes(j)%dofs(k)
565 amg%lvl(1)%map_finest2lvl(nid) = amg%lvl(1)%nodes(j)%gid
566 end do
567 end do
568 ! map_finest2lvl is allocated 0:fine_lvl_dofs (element 0 carries the
569 ! length for the device kernels), so size() is one more than the number
570 ! of dofs and must not be used as the upper loop bound.
571 n = amg%lvl(1)%fine_lvl_dofs
572 do l = 2, amg%nlvls
573 ! Invert the level's node->dof lists once into dof2gid, then compose
574 ! with the level below. Searching every node for every dof instead
575 ! is O(n * dofs-on-level) and dominates setup on large coarse grids.
576 allocate(dof2gid(amg%lvl(l)%fine_lvl_dofs))
577 dof2gid = 0
578 do j = 1, amg%lvl(l)%nnodes
579 do k = 1, amg%lvl(l)%nodes(j)%ndofs
580 dof2gid(amg%lvl(l)%nodes(j)%dofs(k)) = amg%lvl(l)%nodes(j)%gid
581 end do
582 end do
583 do i = 1, n
584 nid = amg%lvl(l-1)%map_finest2lvl(i)
585 if (nid .ge. 1 .and. nid .le. amg%lvl(l)%fine_lvl_dofs) then
586 if (dof2gid(nid) .gt. 0) then
587 amg%lvl(l)%map_finest2lvl(i) = dof2gid(nid)
588 end if
589 end if
590 end do
591 deallocate(dof2gid)
592 end do
593 if (neko_bcknd_device .eq. 1) then
594 ! The whole array, element 0 included, is what gets mirrored; keep
595 ! this independent of the loop bound above.
596 nmap = size(amg%lvl(1)%map_finest2lvl)
597 do l = 1, amg%nlvls
598 amg%lvl(l)%map_finest2lvl(0) = nmap
599 call device_memcpy( amg%lvl(l)%map_finest2lvl, &
600 amg%lvl(l)%map_finest2lvl_d, nmap, &
601 host_to_device, .true.)
602 call device_memcpy( amg%lvl(l)%map_f2c, &
603 amg%lvl(l)%map_f2c_d, amg%lvl(l)%fine_lvl_dofs+1, &
604 host_to_device, .true.)
605 end do
606 end if
607 end subroutine fill_lvl_map
608
615 subroutine build_agg_csr(amg)
616 type(tamg_hierarchy_t), intent(inout) :: amg
620 integer, parameter :: PART_PAD = 32
621 integer, allocatable :: fill(:), seen(:)
622 integer :: l, i, j, k, c, n, nagg, nthrds, ldpart
623
624 n = amg%lvl(1)%fine_lvl_dofs
625
626 nthrds = 1
627 !$ nthrds = omp_get_max_threads()
628
629 ! The restriction and prolongation operators thread over nodes on the
630 ! assumption that a level's nodes partition the dofs of the level below,
631 ! which is what makes their writes disjoint. Check it once here rather
632 ! than let a violation turn into a silent data race at solve time.
633 do l = 1, amg%nlvls
634 allocate(seen(amg%lvl(l)%fine_lvl_dofs))
635 seen = 0
636 do j = 1, amg%lvl(l)%nnodes
637 do k = 1, amg%lvl(l)%nodes(j)%ndofs
638 i = amg%lvl(l)%nodes(j)%dofs(k)
639 if (i .lt. 1 .or. i .gt. amg%lvl(l)%fine_lvl_dofs) then
640 call neko_error('TAMG: node dof outside the level')
641 end if
642 if (seen(i) .ne. 0) then
643 call neko_error('TAMG: dof shared by two nodes on a level')
644 end if
645 seen(i) = j
646 end do
647 end do
648 deallocate(seen)
649 end do
650
651 do l = 1, amg%nlvls
652 nagg = amg%lvl(l)%nnodes
653
654 allocate(amg%lvl(l)%agg_ptr(nagg + 1))
655 allocate(amg%lvl(l)%agg_dof(n))
656 allocate(fill(nagg))
657
658 ! Levels with too few aggregates to give every thread one fall back
659 ! to a reduction over dofs into per-thread partials; the coarsest
660 ! level always does, it has a single aggregate from aggregate_end.
661 ! Pad the leading dimension so every thread's column starts in its
662 ! own cache line: unpadded, nagg = 1 puts all nthrds accumulators in
663 ! a single line and the inner loop degenerates into line ping-pong,
664 ! which is no better than the atomic it replaces. PART_PAD is in
665 ! elements and covers the widest line in play (A64FX is 256 B, x86
666 ! and most AArch64 are 64 B).
667 if (nagg .lt. 2 * nthrds) then
668 ldpart = ((nagg + part_pad - 1) / part_pad) * part_pad
669 allocate(amg%lvl(l)%agg_part(ldpart, nthrds))
670 end if
671
672 associate(aptr => amg%lvl(l)%agg_ptr, adof => amg%lvl(l)%agg_dof, &
673 map => amg%lvl(l)%map_finest2lvl)
674
675 ! Histogram the aggregate sizes into aptr(2:nagg+1). map_finest2lvl
676 ! is never initialised, so a dof left out of the aggregation shows
677 ! up here rather than as silent garbage in the matvec.
678 aptr = 0
679 do i = 1, n
680 c = map(i)
681 if (c .lt. 1 .or. c .gt. nagg) then
682 call neko_error('TAMG: dof not covered by the aggregation')
683 end if
684 aptr(c + 1) = aptr(c + 1) + 1
685 end do
686
687 ! Prefix sum into 1-based offsets
688 aptr(1) = 1
689 do c = 1, nagg
690 aptr(c + 1) = aptr(c + 1) + aptr(c)
691 end do
692
693 ! Place the dof ids. i ascends, so each aggregate's list comes out
694 ! sorted, which keeps the indirect reads in the matvec monotone.
695 fill = 0
696 do i = 1, n
697 c = map(i)
698 adof(aptr(c) + fill(c)) = i
699 fill(c) = fill(c) + 1
700 end do
701 end associate
702
703 deallocate(fill)
704 end do
705
706 end subroutine build_agg_csr
707end module tree_amg_multigrid
__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.
Definition device.F90:113
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 Matrix-vector product.
Definition ax.f90:34
Defines a list of bc_t.
Definition bc_list.f90:34
Coefficients.
Definition coef.f90:34
Definition comm.F90:1
type(mpi_comm), public neko_comm
MPI communicator.
Definition comm.F90:46
subroutine, public device_sub3(a_d, b_d, c_d, n, strm)
Vector subtraction .
subroutine, public device_add2(a_d, b_d, n, strm)
Vector addition .
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
subroutine, public device_col2(a_d, b_d, n, strm)
Vector multiplication .
real(kind=rp) function, public device_glsc2(a_d, b_d, n, strm)
Weighted inner product .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
integer, parameter, public host_to_device
Definition device.F90:48
Gather-scatter.
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
NEKTON map.
Definition map.f90:3
Definition math.f90:60
subroutine, public add2s1(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on first argument)
Definition math.f90:981
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 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 rzero(a, n)
Zero a real vector.
Definition math.f90:235
Defines a mesh.
Definition mesh.f90:34
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
Defines a function space.
Definition space.f90:34
Implements an aggregation for TreeAMG hierarchy structure.
subroutine aggregate_end(tamg, lvl_id)
Aggregate all dofs to a single point to form a tree-like structure.
subroutine aggregate_pairs(tamg, lvl_id, max_aggs, facet_neigh, agg_nhbr)
Aggregates pairs of dofs based on adjacent dofs.
subroutine aggregate_greedy(tamg, lvl_id, max_aggs, facet_neigh, agg_nhbr)
Aggregates dofs based on adjacent dofs.
subroutine aggregate_finest_level(tamg, lx, ly, lz, ne)
Aggregaiton on finest level Aggregates all dofs in an element into a single aggregate.
Implements multigrid using the TreeAMG hierarchy structure. USE:
subroutine calc_resid(r, x, b, amg, lvl, n)
Wrapper function to calculate residyal.
subroutine tamg_mg_set_eig_refresh(this, warm_start, power_its_refresh)
Set the eigenvalue re-estimation policy on every level.
subroutine tamg_mg_cycle_d(this, zero_initial_guess)
multigrid cycle for the TreeAMG solver object on device
subroutine build_agg_csr(amg)
Build the transpose of map_finest2lvl in CSR form: for every aggregate on a level,...
subroutine tamg_mg_invalidate_eigs(this)
Re-estimate every level's eigenvalues on the next solve. Needed when the operator changes underneath ...
subroutine tamg_mg_free(this)
free tree amg solver object
subroutine fill_lvl_map(amg)
Create index mapping between levels and directly to finest level.
subroutine print_resid_info(r, x, b, r_d, x_d, b_d, amg, lvl, n)
subroutine tamg_mg_solve(this, z, r, n)
Solver function for the TreeAMG solver object.
subroutine tamg_mg_cycle(this, zero_initial_guess)
multigrid cycle for the TreeAMG solver object
subroutine print_preagg_info(lvl, nagg, agg_type)
subroutine tamg_mg_init(this, ax, xh, coef, msh, gs_h, nlvls, blst, max_iter, cheby_degree)
Initialization of the TreeAMG multigrid solver.
Implements smoothers for use with TreeAMG matrix vector product.
Implements the base type for TreeAMG hierarchy structure.
Definition tree_amg.f90:34
subroutine, public tamg_node_init(node, gid, ndofs)
Initialization of a TreeAMG tree node.
Definition tree_amg.f90:280
subroutine, public tamg_lvl_init(tamg_lvl, lvl, nnodes, ndofs)
Initialization of a TreeAMG level.
Definition tree_amg.f90:199
Utilities.
Definition utils.f90:35
subroutine, public neko_warning(warning_msg)
Reports a warning to standard output.
Definition utils.f90:398
Base type for a matrix-vector product providing .
Definition ax.f90:43
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:63
Gather-scatter kernel.
The function space for the SEM solution fields.
Definition space.f90:64
Type for a TreeAMG hierarchy.
Definition tree_amg.f90:100
Type for the TreeAMG solver.
Type for Chebyshev iteration using TreeAMG matvec.