Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
phmg.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!
34module phmg
35 use num_types, only : rp
36 use precon, only : pc_t
37 use gather_scatter, only : gs_t, gs_op_add
38 use space, only : space_t, gll
39 use dofmap, only : dofmap_t
40 use field, only : field_t
42 use mesh, only : mesh_t
43 use bc, only : bc_t
44 use bc_list, only : bc_list_t
46 use dirichlet, only : dirichlet_t
47 use utils, only : neko_error, neko_warning
48 use cheby, only : cheby_t
50 use jacobi, only : jacobi_t
52 use schwarz, only : schwarz_t
53 use ax_product, only : ax_t, ax_helm_allocator
56 use json_module, only : json_file
58 use math, only : copy, col2, add2, add2s2, add2s1
64 use krylov, only : ksp_t, ksp_monitor_t, ksp_max_iter, &
65 krylov_solver_factory
67 use logger, only : neko_log, log_size
68 use, intrinsic :: iso_c_binding
69 implicit none
70 private
71
72
73 type, private :: phmg_lvl_t
74 integer :: lvl = -1
75 integer :: smoother_itrs = 10
76 type(space_t), pointer :: xh
77 type(dofmap_t), pointer :: dm_xh
78 type(gs_t), pointer :: gs_h
80 type(cheby_t) :: cheby
84 type(coef_t), pointer :: coef
85 type(scalar_bc_projector_t) :: bc_projector
86 type(dirichlet_t) :: bc
87 type(field_t) :: r, w, z
88 end type phmg_lvl_t
89
90 type, public :: phmg_hrchy_t
91 type(phmg_lvl_t), allocatable :: lvl(:)
92 end type phmg_hrchy_t
93
94
95 type, public, extends(pc_t) :: phmg_t
96 type(tamg_solver_t) :: amg_solver
97 integer :: nlvls
98 type(phmg_hrchy_t) :: phmg_hrchy
99 class(ax_t), allocatable :: ax
100 type(interpolator_t), allocatable :: intrp(:)
102 type(interpolator_t), allocatable :: crd_intrp(:)
103 type(mesh_t), pointer :: msh
105 integer :: last_metrics_version = -1
109 logical :: update_enabled = .false.
111 logical :: refresh_eigs = .true.
113 integer :: refresh_eigs_frequency = 20
115 logical :: eigs_warm_start = .true.
117 integer :: power_its_refresh = 20
119 integer :: n_refresh = 0
121 character(len=:), allocatable :: cheby_acc
122 contains
123 procedure, pass(this) :: init => phmg_init
124 procedure, pass(this) :: init_from_components => &
126 procedure, pass(this) :: free => phmg_free
127 procedure, pass(this) :: solve => phmg_solve
128 procedure, pass(this) :: update => phmg_update
129 procedure, private, pass(this) :: mg_cycle => phmg_mg_cycle
130 end type phmg_t
131
132contains
133
134 subroutine phmg_init(this, coef, bclst, phmg_params)
135 class(phmg_t), intent(inout), target :: this
136 type(coef_t), intent(in), target :: coef
137 type(bc_list_t), intent(inout), target :: bclst
138 type(json_file), intent(inout) :: phmg_params
139 integer :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
140 integer :: smoother_itrs
141 character(len=:), allocatable :: cheby_acc
142 integer, allocatable :: pcrs_sched(:)
143 logical :: update_enabled
144
145 call json_get_or_default(phmg_params, 'smoother_iterations', &
146 smoother_itrs, 3)
147
148 call json_get_or_default(phmg_params, 'smoother_cheby_acc', &
149 cheby_acc, "jacobi")
150
151 call json_get_or_default(phmg_params, 'coarse_grid.levels', &
152 crs_tamg_lvls, 3)
153
154 call json_get_or_default(phmg_params, 'coarse_grid.iterations', &
155 crs_tamg_itrs, 1)
156
157 call json_get_or_default(phmg_params, 'coarse_grid.cheby_degree', &
158 crs_tamg_cheby_degree, 4)
159
160 if (phmg_params%valid_path('pcoarsening_schedule')) then
161 call json_get(phmg_params, 'pcoarsening_schedule', pcrs_sched)
162 else
163 allocate(pcrs_sched(2))
164 pcrs_sched(1) = 3
165 pcrs_sched(2) = 1
166 end if
167
168 call json_get_or_default(phmg_params, 'update.enabled', &
169 update_enabled, .false.)
170
171 ! Control the eigenvalue re-estimation; geometry always refreshes.
172 call json_get_or_default(phmg_params, 'update.eigs.enabled', &
173 this%refresh_eigs, .true.)
174 call json_get_or_default(phmg_params, 'update.eigs.frequency', &
175 this%refresh_eigs_frequency, 20)
176 call json_get_or_default(phmg_params, 'update.eigs.warm_start', &
177 this%eigs_warm_start, .true.)
178 call json_get_or_default(phmg_params, 'update.eigs.warm_start_iterations', &
179 this%power_its_refresh, 20)
180
181 call this%init_from_components(coef, bclst, smoother_itrs, &
182 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree,&
183 pcrs_sched, update_enabled)
184
185 end subroutine phmg_init
186
187 subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, &
188 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, &
189 pcrs_sched, update_enabled)
190 class(phmg_t), intent(inout), target :: this
191 type(coef_t), intent(in), target :: coef
192 type(bc_list_t), intent(inout), target :: bclst
193 integer, intent(in) :: smoother_itrs
194 character(len=:), allocatable :: cheby_acc
195 integer, intent(in) :: crs_tamg_lvls, crs_tamg_itrs
196 integer, intent(in) :: crs_tamg_cheby_degree
197 integer, intent(in), allocatable :: pcrs_sched(:)
198 logical, intent(in), optional :: update_enabled
199 integer :: lx_crs, lx_mid
200 integer, allocatable :: lx_lvls(:)
201 integer :: n, i, j, st
202 class(bc_t), pointer :: bc_j
203 logical :: use_jacobi, use_cheby
204 use_jacobi = .true.
205 use_cheby = .true.
206
207 this%msh => coef%msh
208 this%cheby_acc = trim(cheby_acc)
209
210 if (present(update_enabled)) then
211 this%update_enabled = update_enabled
212 else
213 this%update_enabled = .false.
214 end if
215
216 this%nlvls = size(pcrs_sched) + 1
217 allocate(lx_lvls(0:this%nlvls - 1))
218 lx_lvls(1:) = pcrs_sched + 1
219
220 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
221
222 this%phmg_hrchy%lvl(0)%lvl = 0
223 this%phmg_hrchy%lvl(0)%smoother_itrs = smoother_itrs
224 this%phmg_hrchy%lvl(0)%Xh => coef%Xh
225 this%phmg_hrchy%lvl(0)%coef => coef
226 this%phmg_hrchy%lvl(0)%dm_Xh => coef%dof
227 this%phmg_hrchy%lvl(0)%gs_h => coef%gs_h
228
229 do i = 1, this%nlvls - 1
230 allocate(this%phmg_hrchy%lvl(i)%Xh)
231 allocate(this%phmg_hrchy%lvl(i)%dm_Xh)
232 allocate(this%phmg_hrchy%lvl(i)%gs_h)
233 allocate(this%phmg_hrchy%lvl(i)%coef)
234
235 this%phmg_hrchy%lvl(i)%lvl = i
236 this%phmg_hrchy%lvl(i)%smoother_itrs = smoother_itrs
237 call this%phmg_hrchy%lvl(i)%Xh%init(gll, lx_lvls(i), lx_lvls(i), &
238 lx_lvls(i))
239 call this%phmg_hrchy%lvl(i)%dm_Xh%init(coef%msh, &
240 this%phmg_hrchy%lvl(i)%Xh)
241 call this%phmg_hrchy%lvl(i)%gs_h%init(this%phmg_hrchy%lvl(i)%dm_Xh)
242 ! A coarse level reads G_ij, h1, h2, B and mult and nothing else, so
243 ! retaining the rest costs memory for no purpose. The exception is a
244 ! moving mesh, where update() has to rebuild the geometry from the
245 ! derivative arrays.
246 if (this%update_enabled) then
247 call this%phmg_hrchy%lvl(i)%coef%init( &
248 this%phmg_hrchy%lvl(i)%gs_h, coef_full)
249 else
250 call this%phmg_hrchy%lvl(i)%coef%init( &
251 this%phmg_hrchy%lvl(i)%gs_h, coef_operator)
252 end if
253 end do
254
255 do i = 0, this%nlvls - 1
256 call this%phmg_hrchy%lvl(i)%r%init(this%phmg_hrchy%lvl(i)%dm_Xh)
257 call this%phmg_hrchy%lvl(i)%w%init(this%phmg_hrchy%lvl(i)%dm_Xh)
258 call this%phmg_hrchy%lvl(i)%z%init(this%phmg_hrchy%lvl(i)%dm_Xh)
259
260 this%phmg_hrchy%lvl(i)%coef%ifh2 = coef%ifh2
261 call copy(this%phmg_hrchy%lvl(i)%coef%h1, coef%h1, &
262 this%phmg_hrchy%lvl(i)%dm_Xh%size())
263
264 call this%phmg_hrchy%lvl(i)%bc%init_base(this%phmg_hrchy%lvl(i)%coef)
265 if (bclst%size() .gt. 0 ) then
266 do j = 1, bclst%size()
267 bc_j => bclst%get(j)
268 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
269 end do
270 end if
271 call this%phmg_hrchy%lvl(i)%bc%finalize()
272 call this%phmg_hrchy%lvl(i)%bc%set_g(0.0_rp)
273 call this%phmg_hrchy%lvl(i)%bc_projector%mark(this%phmg_hrchy%lvl(i)%bc)
274
276 if (trim(cheby_acc) .eq. "schwarz") then
277 call this%phmg_hrchy%lvl(i)%schwarz%init( &
278 this%phmg_hrchy%lvl(i)%Xh, &
279 this%phmg_hrchy%lvl(i)%dm_Xh, &
280 this%phmg_hrchy%lvl(i)%gs_h, &
281 this%phmg_hrchy%lvl(i)%bc_projector, &
282 coef%msh)
283 end if
284
285 if (neko_bcknd_device .eq. 1) then
286 call this%phmg_hrchy%lvl(i)%device_jacobi%init(&
287 this%phmg_hrchy%lvl(i)%coef, &
288 this%phmg_hrchy%lvl(i)%dm_Xh, &
289 this%phmg_hrchy%lvl(i)%gs_h)
290 else
291 call this%phmg_hrchy%lvl(i)%jacobi%init(&
292 this%phmg_hrchy%lvl(i)%coef, &
293 this%phmg_hrchy%lvl(i)%dm_Xh, &
294 this%phmg_hrchy%lvl(i)%gs_h)
295 end if
296
297 if (neko_bcknd_device .eq. 1) then
298 if (trim(cheby_acc) .eq. "jacobi") then
299 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
300 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
301 this%phmg_hrchy%lvl(i)%device_jacobi)
302 st = 1
303 else
304 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
305 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
306 st = 0
307 if (trim(cheby_acc) .eq. "schwarz") then
308 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
309 this%phmg_hrchy%lvl(i)%schwarz
310 st = 2
311 end if
312 end if
313 else
314 if (trim(cheby_acc) .eq. "jacobi") then
315 call this%phmg_hrchy%lvl(i)%cheby%init( &
316 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
317 this%phmg_hrchy%lvl(i)%jacobi)
318 st = 1
319 else
320 call this%phmg_hrchy%lvl(i)%cheby%init( &
321 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
322 st = 0
323 if (trim(cheby_acc) .eq. "schwarz") then
324 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
325 this%phmg_hrchy%lvl(i)%schwarz
326 st = 2
327 end if
328 end if
329 end if
330
331 end do
332
333 call print_phmg_info(this%nlvls, st, this%phmg_hrchy)
334
335 ! Create backend specific Ax operator
336 call ax_helm_allocator(this%ax, type_name = "standard")
337
338 ! Interpolator Fine + mg levels
339 allocate(this%intrp(this%nlvls - 1))
340 do i = 1, this%nlvls -1
341 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
342 this%phmg_hrchy%lvl(i)%Xh)
343 end do
344
345 ! Coarse space first. Each level maps from level 0.
346 if (this%update_enabled) then
347 allocate(this%crd_intrp(this%nlvls - 1))
348 do i = 1, this%nlvls - 1
349 call this%crd_intrp(i)%init(this%phmg_hrchy%lvl(i)%Xh, &
350 this%phmg_hrchy%lvl(0)%Xh)
351 end do
352 end if
353
354 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
355 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
356 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
357 this%phmg_hrchy%lvl(this%nlvls -1)%bc_projector, &
358 crs_tamg_itrs, crs_tamg_cheby_degree)
359
360 ! update() only refreshes when `lvl(0)%coef%metrics_version` changes.
361 this%last_metrics_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
362
363 ! Hand the eigenvalue policy to every smoother.
364 if (this%update_enabled) then
365 do i = 0, this%nlvls - 1
366 this%phmg_hrchy%lvl(i)%cheby%warm_start_eigs = this%eigs_warm_start
367 this%phmg_hrchy%lvl(i)%cheby%power_its_refresh = &
368 this%power_its_refresh
369 this%phmg_hrchy%lvl(i)%cheby_device%warm_start_eigs = &
370 this%eigs_warm_start
371 this%phmg_hrchy%lvl(i)%cheby_device%power_its_refresh = &
372 this%power_its_refresh
373 end do
374 call this%amg_solver%set_eig_refresh(this%eigs_warm_start, &
375 this%power_its_refresh)
376
377 if (trim(cheby_acc) .eq. "schwarz") then
378 call neko_warning("PHMG: the Schwarz smoother is not refreshed " // &
379 "when the mesh changes. Its local solves stay at the " // &
380 "initial geometry.")
381 end if
382 end if
383
384 end subroutine phmg_init_from_components
385
386 subroutine phmg_free(this)
387 class(phmg_t), intent(inout) :: this
388 integer :: i
389
390 call this%amg_solver%free()
391
392 if (allocated(this%intrp)) then
393 do i = 1, size(this%intrp)
394 call this%intrp(i)%free()
395 end do
396 deallocate(this%intrp)
397 end if
398
399 if (allocated(this%crd_intrp)) then
400 do i = 1, size(this%crd_intrp)
401 call this%crd_intrp(i)%free()
402 end do
403 deallocate(this%crd_intrp)
404 end if
405
406 if (allocated(this%ax)) then
407 call this%ax%free()
408 deallocate(this%ax)
409 end if
410
411 if (allocated(this%phmg_hrchy%lvl)) then
412 do i = lbound(this%phmg_hrchy%lvl, 1), ubound(this%phmg_hrchy%lvl, 1)
413 call this%phmg_hrchy%lvl(i)%r%free()
414 call this%phmg_hrchy%lvl(i)%w%free()
415 call this%phmg_hrchy%lvl(i)%z%free()
416
417 call this%phmg_hrchy%lvl(i)%cheby%free()
418 call this%phmg_hrchy%lvl(i)%cheby_device%free()
419 call this%phmg_hrchy%lvl(i)%jacobi%free()
420 call this%phmg_hrchy%lvl(i)%device_jacobi%free()
421
422 if (allocated(this%phmg_hrchy%lvl(i)%schwarz%work1)) then
423 call this%phmg_hrchy%lvl(i)%schwarz%free()
424 end if
425
426 call this%phmg_hrchy%lvl(i)%bc_projector%free()
427 call this%phmg_hrchy%lvl(i)%bc%free()
428
429 ! Level 0 borrows Xh, dm_Xh, gs_h and coef from the caller,
430 ! all other levels own them
431 if (this%phmg_hrchy%lvl(i)%lvl .gt. 0) then
432 call this%phmg_hrchy%lvl(i)%coef%free()
433 call this%phmg_hrchy%lvl(i)%gs_h%free()
434 call this%phmg_hrchy%lvl(i)%dm_Xh%free()
435 call this%phmg_hrchy%lvl(i)%Xh%free()
436 deallocate(this%phmg_hrchy%lvl(i)%coef)
437 deallocate(this%phmg_hrchy%lvl(i)%gs_h)
438 deallocate(this%phmg_hrchy%lvl(i)%dm_Xh)
439 deallocate(this%phmg_hrchy%lvl(i)%Xh)
440 end if
441
442 nullify(this%phmg_hrchy%lvl(i)%coef)
443 nullify(this%phmg_hrchy%lvl(i)%gs_h)
444 nullify(this%phmg_hrchy%lvl(i)%dm_Xh)
445 nullify(this%phmg_hrchy%lvl(i)%Xh)
446 end do
447 deallocate(this%phmg_hrchy%lvl)
448 end if
449
450 nullify(this%msh)
451
452 end subroutine phmg_free
453
454 subroutine phmg_solve(this, z, r, n)
455 class(phmg_t), intent(inout) :: this
456 integer, intent(in) :: n
457 real(kind=rp), dimension(n), intent(inout) :: z
458 real(kind=rp), dimension(n), intent(inout) :: r
459 type(c_ptr) :: z_d, r_d
460 type(ksp_monitor_t) :: ksp_results
461 integer :: i
462
463 call profiler_start_region('PHMG_solve', 8)
464 associate( mglvl => this%phmg_hrchy%lvl)
465 if (neko_bcknd_device .eq. 1) then
466 z_d = device_get_ptr(z)
467 r_d = device_get_ptr(r)
468 !We should not work with the input
469 call device_copy(mglvl(0)%r%x_d, r_d, n)
470
471 call device_rzero(mglvl(0)%z%x_d, n)
472 call device_rzero(mglvl(0)%w%x_d, n)
473
474 call this%mg_cycle()
475
476 call device_copy(z_d, mglvl(0)%z%x_d, n)
477 else
478 !OCL NORECURRENCE, NOVREC, NOALIAS
479 !DIR$ CONCURRENT
480 !DIR$ IVDEP
481 !GCC$ ivdep
482 !$omp parallel do
483 do i = 1, n
484 !We should not work with the input
485 mglvl(0)%r%x(i,1,1,1) = r(i)
486
487 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
488 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
489 end do
490 !$omp end parallel do
491
492 call this%mg_cycle()
493
494 !OCL NORECURRENCE, NOVREC, NOALIAS
495 !DIR$ CONCURRENT
496 !DIR$ IVDEP
497 !GCC$ ivdep
498 !$omp parallel do
499 do i = 1, n
500 z(i) = mglvl(0)%z%x(i,1,1,1)
501 end do
502 !$omp end parallel do
503 end if
504 end associate
505 call profiler_end_region('PHMG_solve', 8)
506
507 end subroutine phmg_solve
508
510 subroutine phmg_update(this)
511 class(phmg_t), intent(inout) :: this
512 integer :: fine_version
513 logical :: do_eigs
514
515 if (.not. this%update_enabled) return
516
517 fine_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
518
519 ! Mesh has not changed since the last refresh.
520 if (fine_version .eq. this%last_metrics_version) return
521
523
524 call phmg_update_smoother_acc(this)
525
526 do_eigs = (this%refresh_eigs) .and. &
527 (this%refresh_eigs_frequency .gt. 0) .and. &
528 (mod(this%n_refresh, max(this%refresh_eigs_frequency, 1)) &
529 .eq. 0)
530 if (do_eigs) call phmg_update_smoother_eigs(this)
531
532 this%n_refresh = this%n_refresh + 1
533 this%last_metrics_version = fine_version
534
535 end subroutine phmg_update
536
537
541 class(phmg_t), intent(inout) :: this
542 integer :: i
543
544 if (.not. allocated(this%crd_intrp)) then
545 call neko_error("PHMG: update requested but coordinate " // &
546 "interpolators were not initialized.")
547 end if
548
549 call profiler_start_region('PHMG_update_geometry')
550 associate(mg => this%phmg_hrchy%lvl, nelv => this%msh%nelv)
551 do i = 1, this%nlvls - 1
552 call this%crd_intrp(i)%map(mg(i)%dm_Xh%x%x, mg(0)%dm_Xh%x%x, &
553 nelv, mg(i)%Xh)
554 call this%crd_intrp(i)%map(mg(i)%dm_Xh%y%x, mg(0)%dm_Xh%y%x, &
555 nelv, mg(i)%Xh)
556 call this%crd_intrp(i)%map(mg(i)%dm_Xh%z%x, mg(0)%dm_Xh%z%x, &
557 nelv, mg(i)%Xh)
558
559 call mg(i)%coef%recompute_metrics()
560 end do
561 end associate
562 call profiler_end_region('PHMG_update_geometry')
563
564 end subroutine phmg_update_coarse_geometry
565
566
570 class(phmg_t), intent(inout) :: this
571 integer :: i
572
573 call profiler_start_region('PHMG_update_smoother_acc')
574 associate(mg => this%phmg_hrchy%lvl)
575 select case (trim(this%cheby_acc))
576 case ("jacobi")
577 do i = 0, this%nlvls - 1
578 if (neko_bcknd_device .eq. 1) then
579 call mg(i)%device_jacobi%update()
580 else
581 call mg(i)%jacobi%update()
582 end if
583 end do
584 case ("schwarz")
585 ! Not refreshed. The local solves are built by the FDM from the
586 ! element coordinates, and fdm_t has no update path.
587 case default
588 ! No accelerator to refresh.
589 end select
590 end associate
591 call profiler_end_region('PHMG_update_smoother_acc')
592
593 end subroutine phmg_update_smoother_acc
594
595
598 class(phmg_t), intent(inout) :: this
599 integer :: i
600
601 associate(mg => this%phmg_hrchy%lvl)
602 do i = 0, this%nlvls - 1
603 if (neko_bcknd_device .eq. 1) then
604 mg(i)%cheby_device%recompute_eigs = .true.
605 else
606 mg(i)%cheby%recompute_eigs = .true.
607 end if
608 end do
609 end associate
610
611 call this%amg_solver%invalidate_eigs()
612
613 end subroutine phmg_update_smoother_eigs
614
615
616 subroutine phmg_mg_cycle(this)
617 class(phmg_t), intent(inout) :: this
618 type(ksp_monitor_t) :: ksp_results
619 character(len=2) :: lvl_name
620 integer :: lvl, i
621
622 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
623 msh => this%msh, ax => this%Ax)
624 do lvl = 0, this%nlvls-2
625 write(lvl_name, '(I0)') lvl
626 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
627 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
628 !------------!
629 ! SMOOTH !
630 !------------!
631 if (neko_bcknd_device .eq. 1) then
632 mg(lvl)%cheby_device%zero_initial_guess = .true.
633 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
634 r%x, mg(lvl)%dm_Xh%size(), &
635 mg(lvl)%coef, mg(lvl)%bc_projector, &
636 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
637 else
638 mg(lvl)%cheby%zero_initial_guess = .true.
639 ksp_results = mg(lvl)%cheby%solve(ax, z, &
640 r%x, mg(lvl)%dm_Xh%size(), &
641 mg(lvl)%coef, mg(lvl)%bc_projector, &
642 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
643 end if
644
645 !------------!
646 ! Residual !
647 !------------!
648 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
649 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
650 glb_cmd_event)
651 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
652 call mg(lvl)%bc_projector%apply(w%x, mg(lvl)%dm_Xh%size())
653
654 if (neko_bcknd_device .eq. 1) then
655 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
656 else
657 !OCL NORECURRENCE, NOVREC, NOALIAS
658 !DIR$ CONCURRENT
659 !DIR$ IVDEP
660 !GCC$ ivdep
661 !$omp parallel do
662 do i = 1, mg(lvl)%dm_Xh%size()
663 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
664 end do
665 !$omp end parallel do
666 end if
667
668 !------------!
669 ! Restrict !
670 !------------!
671 if (neko_bcknd_device .eq. 1) then
672 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
673 else
674 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
675 end if
676
677 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
678
679 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
680 gs_op_add, glb_cmd_event)
681 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
682
683 call mg(lvl+1)%bc_projector%apply( &
684 mg(lvl+1)%r%x, &
685 mg(lvl+1)%dm_Xh%size())
686
687 if (neko_bcknd_device .eq. 1) then
688 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
689 else
690 !OCL NORECURRENCE, NOVREC, NOALIAS
691 !DIR$ CONCURRENT
692 !DIR$ IVDEP
693 !GCC$ ivdep
694 !$omp parallel do
695 do i = 1, mg(lvl+1)%dm_Xh%size()
696 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
697 end do
698 !$omp end parallel do
699 end if
700 end associate
701 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
702 end do
703
704 call profiler_start_region( 'PHMG_coarse-solve' )
705 !------------!
706 ! SOLVE !
707 !------------!
708 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
709 mg(this%nlvls-1)%r%x, &
710 mg(this%nlvls-1)%dm_Xh%size())
711 call profiler_end_region( 'PHMG_coarse-solve' )
712
713 do lvl = (this%nlvls-2), 0, -1
714 write(lvl_name, '(I0)') lvl
715 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
716 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
717 !------------!
718 ! Project !
719 !------------!
720 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
721
722 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
723 glb_cmd_event)
724 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
725
726 if (neko_bcknd_device .eq. 1) then
727 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
728 else
729 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
730 end if
731
732 !------------!
733 ! Correct !
734 !------------!
735 if (neko_bcknd_device .eq. 1) then
736 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
737 else
738 !OCL NORECURRENCE, NOVREC, NOALIAS
739 !DIR$ CONCURRENT
740 !DIR$ IVDEP
741 !GCC$ ivdep
742 !$omp parallel do
743 do i = 1, mg(lvl)%dm_Xh%size()
744 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
745 end do
746 !$omp end parallel do
747 end if
748
749 !------------!
750 ! SMOOTH !
751 !------------!
752 if (neko_bcknd_device .eq. 1) then
753 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
754 r%x, mg(lvl)%dm_Xh%size(), &
755 mg(lvl)%coef, mg(lvl)%bc_projector, &
756 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
757 else
758 ksp_results = mg(lvl)%cheby%solve(ax, z, &
759 r%x, mg(lvl)%dm_Xh%size(), &
760 mg(lvl)%coef, mg(lvl)%bc_projector, &
761 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
762 end if
763 end associate
764 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
765 end do
766 end associate
767
768 end subroutine phmg_mg_cycle
769
779 subroutine phmg_jacobi_smoother(z, r, w, mg, msh, Ax, n, lvl)
780 type(phmg_lvl_t) :: mg
781 class(ax_t), intent(inout) :: Ax
782 type(mesh_t), intent(inout) :: msh
783 type(field_t), intent(inout) :: z, r, w
784 integer, intent(in) :: n, lvl
785 integer :: i, j, iblk, ni, niblk
786
787 ni = mg%smoother_itrs
788 if (neko_bcknd_device .eq. 1) then
789 do i = 1, ni
790 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
791 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
792 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
793 call mg%bc_projector%apply(w%x, n)
794 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
795
796 call mg%device_jacobi%solve(w%x, w%x, n)
797
798 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
799 end do
800 else
801 do i = 1, ni
802 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
803 call mg%gs_h%op(w%x, n, gs_op_add)
804 call mg%bc_projector%apply(w%x, n)
805 call add2s1(w%x, r%x, -1.0_rp, n)
806
807 call mg%jacobi%solve(w%x, w%x, n)
808
809 call add2s2(z%x, w%x, 0.6_rp, n)
810 end do
811 end if
812 end subroutine phmg_jacobi_smoother
813
814
815 subroutine phmg_resid_monitor(z, r, w, mg, msh, Ax, lvl, typ)
816 integer :: lvl, typ
817 type(phmg_lvl_t) :: mg
818 class(ax_t), intent(inout) :: Ax
819 type(mesh_t), intent(inout) :: msh
820 type(field_t) :: z, r, w
821 real(kind=rp) :: val
822 character(len=LOG_SIZE) :: log_buf
823 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
824 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
825 call mg%bc_projector%apply(w%x, mg%dm_Xh%size())
826 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
827 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
828 if (typ .eq. 1) then
829 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO - PRE', lvl, val
830 else if (typ .eq. 2) then
831 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO -POST', lvl, val
832 else if (typ .eq. 3) then
833 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO- PRE', lvl, val
834 else if (typ .eq. 4) then
835 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO-POST', lvl, val
836 else if (typ .eq. 5) then
837 write(log_buf, '(A15,I4,F12.6)') 'TAMG - PRE', lvl, val
838 else if (typ .eq. 6) then
839 write(log_buf, '(A15,I4,F12.6)') 'TAMG -POST', lvl, val
840 else
841 write(log_buf, '(A15,I4,F12.6)') 'RESID', lvl, val
842 end if
843 call neko_log%message(log_buf)
844 end subroutine phmg_resid_monitor
845
846 subroutine print_phmg_info(nlvls, smoo_type, phmg)
847 integer, intent(in) :: nlvls
848 integer, intent(in) :: smoo_type
849 type(phmg_hrchy_t) :: phmg
850 integer :: i, clvl
851 character(len=LOG_SIZE) :: log_buf, smoo_name
852
853 call neko_log%section('PHMG')
854
855 if (smoo_type .eq. 1) then
856 write(smoo_name, '(A16)') 'CHEBY-acc JACOBI'
857 else if (smoo_type .eq. 2) then
858 write(smoo_name, '(A17)') 'CHEBY-acc SCHWARZ'
859 else
860 write(smoo_name, '(A5)') 'CHEBY'
861 end if
862
863 write(log_buf, '(A28,I2,A8)') &
864 'Creating PHMG hierarchy with', &
865 nlvls, 'levels.'
866 call neko_log%message(log_buf)
867
868 clvl = nlvls - 1
869 do i = 0, nlvls-1
870 write(log_buf, '(A8,I2,A8,I2)') &
871 '-- level', i, '-- lx:', phmg%lvl(i)%Xh%lx
872 call neko_log%message(log_buf)
873
874 if (i .eq. clvl) then
875 write(log_buf, '(A19,A20)') &
876 'Solve:', 'tAMG'
877 call neko_log%message(log_buf)
878 else
879 write(log_buf, '(A22,A20)') &
880 'Smoother:', &
881 trim(smoo_name)
882 call neko_log%message(log_buf)
883
884 write(log_buf, '(A28,I2)') &
885 'Smoother Iters:', &
886 phmg%lvl(i)%smoother_itrs
887 call neko_log%message(log_buf)
888 end if
889 end do
890
891 call neko_log%end_section()
892
893 end subroutine print_phmg_info
894
895end module phmg
__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
Retrieves a parameter by name or assigns a provided default value. In the latter case also adds the m...
Retrieves a parameter by name or throws an error.
Defines a Matrix-vector product.
Definition ax.f90:34
Defines a list of bc_t.
Definition bc_list.f90:34
Defines a boundary condition.
Definition bc.f90:34
Chebyshev preconditioner.
Chebyshev preconditioner.
Definition cheby.f90:34
Coefficients.
Definition coef.f90:34
integer, parameter, public coef_full
Retain every coefficient. The default, and the only scope that supports recompute_metrics(),...
Definition coef.f90:118
integer, parameter, public coef_operator
Retain only what applying a discrete operator needs: , h1, h2, B and mult.
Definition coef.f90:131
Jacobi preconditioner accelerator backend.
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_add2s2(a_d, b_d, c1, n, strm)
Vector addition with scalar multiplication (multiplication on first argument)
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 .
subroutine, public device_invcol2(a_d, b_d, n, strm)
Vector division .
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
subroutine, public device_stream_wait_event(stream, event, flags)
Synchronize a device stream with an event.
Definition device.F90:1544
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
Definition device.F90:52
type(c_ptr), bind(C), public glb_cmd_event
Event for the global command queue.
Definition device.F90:63
Defines a dirichlet boundary condition.
Definition dirichlet.f90:34
Defines a mapping of the degrees of freedom.
Definition dofmap.f90:35
Defines a field.
Definition field.f90:34
Gather-scatter.
Routines to interpolate between different spaces.
Jacobi preconditioner.
Definition pc_jacobi.f90:34
Utilities for retrieving parameters from the case files.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
integer, parameter, public ksp_max_iter
Maximum number of iters.
Definition krylov.f90:52
Logging routines.
Definition log.f90:34
type(log_t), public neko_log
Global log stream.
Definition log.f90:91
integer, parameter, public log_size
Definition log.f90:46
Definition math.f90:60
subroutine, public add2s1(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on first argument)
Definition math.f90:1020
subroutine, public add2(a, b, n)
Vector addition .
Definition math.f90:939
subroutine, public col2(a, b, n)
Vector multiplication .
Definition math.f90:1085
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:295
subroutine, public add2s2(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on second argument)
Definition math.f90:1037
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:14
Hybrid ph-multigrid preconditioner.
Definition phmg.f90:34
subroutine phmg_update_smoother_acc(this)
Rebuild the accelerator the Chebyshev smoother uses, which depends on the geometry.
Definition phmg.f90:570
subroutine phmg_jacobi_smoother(z, r, w, mg, msh, ax, n, lvl)
Wraps jacobi solve as a residual update relaxation method.
Definition phmg.f90:780
subroutine phmg_solve(this, z, r, n)
Definition phmg.f90:455
subroutine phmg_init(this, coef, bclst, phmg_params)
Definition phmg.f90:135
subroutine phmg_update(this)
Bring the preconditioner back in sync after the mesh has changed.
Definition phmg.f90:511
subroutine phmg_update_smoother_eigs(this)
Mark every smoother to re-estimate its eigenvalues on the next solve.
Definition phmg.f90:598
subroutine phmg_update_coarse_geometry(this)
Sample the new fine coordinates on the coarse levels and rebuild their metrics. Level 0 shares the ca...
Definition phmg.f90:541
subroutine phmg_free(this)
Definition phmg.f90:387
subroutine phmg_mg_cycle(this)
Definition phmg.f90:617
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
Definition phmg.f90:816
subroutine print_phmg_info(nlvls, smoo_type, phmg)
Definition phmg.f90:847
subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, pcrs_sched, update_enabled)
Definition phmg.f90:190
Krylov preconditioner.
Definition precon.f90:34
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 scalar_projector_t.
Overlapping schwarz solves.
Definition schwarz.f90:61
Defines a function space.
Definition space.f90:34
integer, parameter, public gll
Definition space.f90:50
Implements multigrid using the TreeAMG hierarchy structure. USE:
Utilities.
Definition utils.f90:35
subroutine, public neko_warning(warning_msg)
Reports a warning to standard output.
Definition utils.f90:452
Base type for a matrix-vector product providing .
Definition ax.f90:43
Base type for a boundary condition.
Definition bc.f90:73
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Defines a Chebyshev preconditioner.
Definition cheby.f90:55
Defines a Chebyshev preconditioner.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:135
Defines a jacobi preconditioner.
Generic Dirichlet boundary condition on .
Definition dirichlet.f90:49
Gather-scatter kernel.
Interpolation between two space::space_t.
Defines a jacobi preconditioner.
Definition pc_jacobi.f90:45
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:57
Base abstract type for a canonical Krylov method, solving .
Definition krylov.f90:74
Defines a canonical Krylov preconditioner.
Definition precon.f90:40
Projector for scalar boundary conditions.
The function space for the SEM solution fields.
Definition space.f90:64
Type for the TreeAMG solver.
#define max(a, b)
Definition tensor.cu:40