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 deallocate(this%ax)
408 end if
409
410 if (allocated(this%phmg_hrchy%lvl)) then
411 do i = lbound(this%phmg_hrchy%lvl, 1), ubound(this%phmg_hrchy%lvl, 1)
412 call this%phmg_hrchy%lvl(i)%r%free()
413 call this%phmg_hrchy%lvl(i)%w%free()
414 call this%phmg_hrchy%lvl(i)%z%free()
415
416 call this%phmg_hrchy%lvl(i)%cheby%free()
417 call this%phmg_hrchy%lvl(i)%cheby_device%free()
418 call this%phmg_hrchy%lvl(i)%jacobi%free()
419 call this%phmg_hrchy%lvl(i)%device_jacobi%free()
420
421 if (allocated(this%phmg_hrchy%lvl(i)%schwarz%work1)) then
422 call this%phmg_hrchy%lvl(i)%schwarz%free()
423 end if
424
425 call this%phmg_hrchy%lvl(i)%bc_projector%free()
426 call this%phmg_hrchy%lvl(i)%bc%free()
427
428 ! Level 0 borrows Xh, dm_Xh, gs_h and coef from the caller,
429 ! all other levels own them
430 if (this%phmg_hrchy%lvl(i)%lvl .gt. 0) then
431 call this%phmg_hrchy%lvl(i)%coef%free()
432 call this%phmg_hrchy%lvl(i)%gs_h%free()
433 call this%phmg_hrchy%lvl(i)%dm_Xh%free()
434 call this%phmg_hrchy%lvl(i)%Xh%free()
435 deallocate(this%phmg_hrchy%lvl(i)%coef)
436 deallocate(this%phmg_hrchy%lvl(i)%gs_h)
437 deallocate(this%phmg_hrchy%lvl(i)%dm_Xh)
438 deallocate(this%phmg_hrchy%lvl(i)%Xh)
439 end if
440
441 nullify(this%phmg_hrchy%lvl(i)%coef)
442 nullify(this%phmg_hrchy%lvl(i)%gs_h)
443 nullify(this%phmg_hrchy%lvl(i)%dm_Xh)
444 nullify(this%phmg_hrchy%lvl(i)%Xh)
445 end do
446 deallocate(this%phmg_hrchy%lvl)
447 end if
448
449 nullify(this%msh)
450
451 end subroutine phmg_free
452
453 subroutine phmg_solve(this, z, r, n)
454 class(phmg_t), intent(inout) :: this
455 integer, intent(in) :: n
456 real(kind=rp), dimension(n), intent(inout) :: z
457 real(kind=rp), dimension(n), intent(inout) :: r
458 type(c_ptr) :: z_d, r_d
459 type(ksp_monitor_t) :: ksp_results
460 integer :: i
461
462 call profiler_start_region('PHMG_solve', 8)
463 associate( mglvl => this%phmg_hrchy%lvl)
464 if (neko_bcknd_device .eq. 1) then
465 z_d = device_get_ptr(z)
466 r_d = device_get_ptr(r)
467 !We should not work with the input
468 call device_copy(mglvl(0)%r%x_d, r_d, n)
469
470 call device_rzero(mglvl(0)%z%x_d, n)
471 call device_rzero(mglvl(0)%w%x_d, n)
472
473 call this%mg_cycle()
474
475 call device_copy(z_d, mglvl(0)%z%x_d, n)
476 else
477 !OCL NORECURRENCE, NOVREC, NOALIAS
478 !DIR$ CONCURRENT
479 !DIR$ IVDEP
480 !GCC$ ivdep
481 !$omp parallel do
482 do i = 1, n
483 !We should not work with the input
484 mglvl(0)%r%x(i,1,1,1) = r(i)
485
486 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
487 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
488 end do
489 !$omp end parallel do
490
491 call this%mg_cycle()
492
493 !OCL NORECURRENCE, NOVREC, NOALIAS
494 !DIR$ CONCURRENT
495 !DIR$ IVDEP
496 !GCC$ ivdep
497 !$omp parallel do
498 do i = 1, n
499 z(i) = mglvl(0)%z%x(i,1,1,1)
500 end do
501 !$omp end parallel do
502 end if
503 end associate
504 call profiler_end_region('PHMG_solve', 8)
505
506 end subroutine phmg_solve
507
509 subroutine phmg_update(this)
510 class(phmg_t), intent(inout) :: this
511 integer :: fine_version
512 logical :: do_eigs
513
514 if (.not. this%update_enabled) return
515
516 fine_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
517
518 ! Mesh has not changed since the last refresh.
519 if (fine_version .eq. this%last_metrics_version) return
520
522
523 call phmg_update_smoother_acc(this)
524
525 do_eigs = (this%refresh_eigs) .and. &
526 (this%refresh_eigs_frequency .gt. 0) .and. &
527 (mod(this%n_refresh, max(this%refresh_eigs_frequency, 1)) &
528 .eq. 0)
529 if (do_eigs) call phmg_update_smoother_eigs(this)
530
531 this%n_refresh = this%n_refresh + 1
532 this%last_metrics_version = fine_version
533
534 end subroutine phmg_update
535
536
540 class(phmg_t), intent(inout) :: this
541 integer :: i
542
543 if (.not. allocated(this%crd_intrp)) then
544 call neko_error("PHMG: update requested but coordinate " // &
545 "interpolators were not initialized.")
546 end if
547
548 call profiler_start_region('PHMG_update_geometry')
549 associate(mg => this%phmg_hrchy%lvl, nelv => this%msh%nelv)
550 do i = 1, this%nlvls - 1
551 call this%crd_intrp(i)%map(mg(i)%dm_Xh%x, mg(0)%dm_Xh%x, &
552 nelv, mg(i)%Xh)
553 call this%crd_intrp(i)%map(mg(i)%dm_Xh%y, mg(0)%dm_Xh%y, &
554 nelv, mg(i)%Xh)
555 call this%crd_intrp(i)%map(mg(i)%dm_Xh%z, mg(0)%dm_Xh%z, &
556 nelv, mg(i)%Xh)
557
558 call mg(i)%coef%recompute_metrics()
559 end do
560 end associate
561 call profiler_end_region('PHMG_update_geometry')
562
563 end subroutine phmg_update_coarse_geometry
564
565
569 class(phmg_t), intent(inout) :: this
570 integer :: i
571
572 call profiler_start_region('PHMG_update_smoother_acc')
573 associate(mg => this%phmg_hrchy%lvl)
574 select case (trim(this%cheby_acc))
575 case ("jacobi")
576 do i = 0, this%nlvls - 1
577 if (neko_bcknd_device .eq. 1) then
578 call mg(i)%device_jacobi%update()
579 else
580 call mg(i)%jacobi%update()
581 end if
582 end do
583 case ("schwarz")
584 ! Not refreshed. The local solves are built by the FDM from the
585 ! element coordinates, and fdm_t has no update path.
586 case default
587 ! No accelerator to refresh.
588 end select
589 end associate
590 call profiler_end_region('PHMG_update_smoother_acc')
591
592 end subroutine phmg_update_smoother_acc
593
594
597 class(phmg_t), intent(inout) :: this
598 integer :: i
599
600 associate(mg => this%phmg_hrchy%lvl)
601 do i = 0, this%nlvls - 1
602 if (neko_bcknd_device .eq. 1) then
603 mg(i)%cheby_device%recompute_eigs = .true.
604 else
605 mg(i)%cheby%recompute_eigs = .true.
606 end if
607 end do
608 end associate
609
610 call this%amg_solver%invalidate_eigs()
611
612 end subroutine phmg_update_smoother_eigs
613
614
615 subroutine phmg_mg_cycle(this)
616 class(phmg_t), intent(inout) :: this
617 type(ksp_monitor_t) :: ksp_results
618 character(len=2) :: lvl_name
619 integer :: lvl, i
620
621 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
622 msh => this%msh, ax => this%Ax)
623 do lvl = 0, this%nlvls-2
624 write(lvl_name, '(I0)') lvl
625 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
626 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
627 !------------!
628 ! SMOOTH !
629 !------------!
630 if (neko_bcknd_device .eq. 1) then
631 mg(lvl)%cheby_device%zero_initial_guess = .true.
632 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
633 r%x, mg(lvl)%dm_Xh%size(), &
634 mg(lvl)%coef, mg(lvl)%bc_projector, &
635 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
636 else
637 mg(lvl)%cheby%zero_initial_guess = .true.
638 ksp_results = mg(lvl)%cheby%solve(ax, z, &
639 r%x, mg(lvl)%dm_Xh%size(), &
640 mg(lvl)%coef, mg(lvl)%bc_projector, &
641 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
642 end if
643
644 !------------!
645 ! Residual !
646 !------------!
647 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
648 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
649 glb_cmd_event)
650 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
651 call mg(lvl)%bc_projector%apply(w%x, mg(lvl)%dm_Xh%size())
652
653 if (neko_bcknd_device .eq. 1) then
654 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
655 else
656 !OCL NORECURRENCE, NOVREC, NOALIAS
657 !DIR$ CONCURRENT
658 !DIR$ IVDEP
659 !GCC$ ivdep
660 !$omp parallel do
661 do i = 1, mg(lvl)%dm_Xh%size()
662 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
663 end do
664 !$omp end parallel do
665 end if
666
667 !------------!
668 ! Restrict !
669 !------------!
670 if (neko_bcknd_device .eq. 1) then
671 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
672 else
673 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
674 end if
675
676 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
677
678 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
679 gs_op_add, glb_cmd_event)
680 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
681
682 call mg(lvl+1)%bc_projector%apply( &
683 mg(lvl+1)%r%x, &
684 mg(lvl+1)%dm_Xh%size())
685
686 if (neko_bcknd_device .eq. 1) then
687 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
688 else
689 !OCL NORECURRENCE, NOVREC, NOALIAS
690 !DIR$ CONCURRENT
691 !DIR$ IVDEP
692 !GCC$ ivdep
693 !$omp parallel do
694 do i = 1, mg(lvl+1)%dm_Xh%size()
695 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
696 end do
697 !$omp end parallel do
698 end if
699 end associate
700 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
701 end do
702
703 call profiler_start_region( 'PHMG_coarse-solve' )
704 !------------!
705 ! SOLVE !
706 !------------!
707 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
708 mg(this%nlvls-1)%r%x, &
709 mg(this%nlvls-1)%dm_Xh%size())
710 call profiler_end_region( 'PHMG_coarse-solve' )
711
712 do lvl = (this%nlvls-2), 0, -1
713 write(lvl_name, '(I0)') lvl
714 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
715 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
716 !------------!
717 ! Project !
718 !------------!
719 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
720
721 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
722 glb_cmd_event)
723 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
724
725 if (neko_bcknd_device .eq. 1) then
726 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
727 else
728 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
729 end if
730
731 !------------!
732 ! Correct !
733 !------------!
734 if (neko_bcknd_device .eq. 1) then
735 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
736 else
737 !OCL NORECURRENCE, NOVREC, NOALIAS
738 !DIR$ CONCURRENT
739 !DIR$ IVDEP
740 !GCC$ ivdep
741 !$omp parallel do
742 do i = 1, mg(lvl)%dm_Xh%size()
743 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
744 end do
745 !$omp end parallel do
746 end if
747
748 !------------!
749 ! SMOOTH !
750 !------------!
751 if (neko_bcknd_device .eq. 1) then
752 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
753 r%x, mg(lvl)%dm_Xh%size(), &
754 mg(lvl)%coef, mg(lvl)%bc_projector, &
755 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
756 else
757 ksp_results = mg(lvl)%cheby%solve(ax, z, &
758 r%x, mg(lvl)%dm_Xh%size(), &
759 mg(lvl)%coef, mg(lvl)%bc_projector, &
760 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
761 end if
762 end associate
763 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
764 end do
765 end associate
766
767 end subroutine phmg_mg_cycle
768
778 subroutine phmg_jacobi_smoother(z, r, w, mg, msh, Ax, n, lvl)
779 type(phmg_lvl_t) :: mg
780 class(ax_t), intent(inout) :: Ax
781 type(mesh_t), intent(inout) :: msh
782 type(field_t), intent(inout) :: z, r, w
783 integer, intent(in) :: n, lvl
784 integer :: i, j, iblk, ni, niblk
785
786 ni = mg%smoother_itrs
787 if (neko_bcknd_device .eq. 1) then
788 do i = 1, ni
789 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
790 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
791 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
792 call mg%bc_projector%apply(w%x, n)
793 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
794
795 call mg%device_jacobi%solve(w%x, w%x, n)
796
797 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
798 end do
799 else
800 do i = 1, ni
801 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
802 call mg%gs_h%op(w%x, n, gs_op_add)
803 call mg%bc_projector%apply(w%x, n)
804 call add2s1(w%x, r%x, -1.0_rp, n)
805
806 call mg%jacobi%solve(w%x, w%x, n)
807
808 call add2s2(z%x, w%x, 0.6_rp, n)
809 end do
810 end if
811 end subroutine phmg_jacobi_smoother
812
813
814 subroutine phmg_resid_monitor(z, r, w, mg, msh, Ax, lvl, typ)
815 integer :: lvl, typ
816 type(phmg_lvl_t) :: mg
817 class(ax_t), intent(inout) :: Ax
818 type(mesh_t), intent(inout) :: msh
819 type(field_t) :: z, r, w
820 real(kind=rp) :: val
821 character(len=LOG_SIZE) :: log_buf
822 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
823 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
824 call mg%bc_projector%apply(w%x, mg%dm_Xh%size())
825 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
826 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
827 if (typ .eq. 1) then
828 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO - PRE', lvl, val
829 else if (typ .eq. 2) then
830 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO -POST', lvl, val
831 else if (typ .eq. 3) then
832 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO- PRE', lvl, val
833 else if (typ .eq. 4) then
834 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO-POST', lvl, val
835 else if (typ .eq. 5) then
836 write(log_buf, '(A15,I4,F12.6)') 'TAMG - PRE', lvl, val
837 else if (typ .eq. 6) then
838 write(log_buf, '(A15,I4,F12.6)') 'TAMG -POST', lvl, val
839 else
840 write(log_buf, '(A15,I4,F12.6)') 'RESID', lvl, val
841 end if
842 call neko_log%message(log_buf)
843 end subroutine phmg_resid_monitor
844
845 subroutine print_phmg_info(nlvls, smoo_type, phmg)
846 integer, intent(in) :: nlvls
847 integer, intent(in) :: smoo_type
848 type(phmg_hrchy_t) :: phmg
849 integer :: i, clvl
850 character(len=LOG_SIZE) :: log_buf, smoo_name
851
852 call neko_log%section('PHMG')
853
854 if (smoo_type .eq. 1) then
855 write(smoo_name, '(A16)') 'CHEBY-acc JACOBI'
856 else if (smoo_type .eq. 2) then
857 write(smoo_name, '(A17)') 'CHEBY-acc SCHWARZ'
858 else
859 write(smoo_name, '(A5)') 'CHEBY'
860 end if
861
862 write(log_buf, '(A28,I2,A8)') &
863 'Creating PHMG hierarchy with', &
864 nlvls, 'levels.'
865 call neko_log%message(log_buf)
866
867 clvl = nlvls - 1
868 do i = 0, nlvls-1
869 write(log_buf, '(A8,I2,A8,I2)') &
870 '-- level', i, '-- lx:', phmg%lvl(i)%Xh%lx
871 call neko_log%message(log_buf)
872
873 if (i .eq. clvl) then
874 write(log_buf, '(A19,A20)') &
875 'Solve:', 'tAMG'
876 call neko_log%message(log_buf)
877 else
878 write(log_buf, '(A22,A20)') &
879 'Smoother:', &
880 trim(smoo_name)
881 call neko_log%message(log_buf)
882
883 write(log_buf, '(A28,I2)') &
884 'Smoother Iters:', &
885 phmg%lvl(i)%smoother_itrs
886 call neko_log%message(log_buf)
887 end if
888 end do
889
890 call neko_log%end_section()
891
892 end subroutine print_phmg_info
893
894end 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:76
integer, parameter, public coef_operator
Retain only what applying a discrete operator needs: , h1, h2, B and mult.
Definition coef.f90:89
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:80
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:984
subroutine, public add2(a, b, n)
Vector addition .
Definition math.f90:903
subroutine, public col2(a, b, n)
Vector multiplication .
Definition math.f90:1049
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:294
subroutine, public add2s2(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on second argument)
Definition math.f90:1001
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:569
subroutine phmg_jacobi_smoother(z, r, w, mg, msh, ax, n, lvl)
Wraps jacobi solve as a residual update relaxation method.
Definition phmg.f90:779
subroutine phmg_solve(this, z, r, n)
Definition phmg.f90:454
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:510
subroutine phmg_update_smoother_eigs(this)
Mark every smoother to re-estimate its eigenvalues on the next solve.
Definition phmg.f90:597
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:540
subroutine phmg_free(this)
Definition phmg.f90:387
subroutine phmg_mg_cycle(this)
Definition phmg.f90:616
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
Definition phmg.f90:815
subroutine print_phmg_info(nlvls, smoo_type, phmg)
Definition phmg.f90:846
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:398
Base type for a matrix-vector product providing .
Definition ax.f90:43
Base type for a boundary condition.
Definition bc.f90:72
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:93
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