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, json)
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) :: json
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(json, 'smoother_iterations', &
146 smoother_itrs, 3)
147
148 call json_get_or_default(json, 'smoother_cheby_acc', &
149 cheby_acc, "jacobi")
150
151 call json_get_or_default(json, 'coarse_grid.levels', &
152 crs_tamg_lvls, 3)
153
154 call json_get_or_default(json, 'coarse_grid.iterations', &
155 crs_tamg_itrs, 1)
156
157 call json_get_or_default(json, 'coarse_grid.cheby_degree', &
158 crs_tamg_cheby_degree, 4)
159
160 if (json%valid_path('pcoarsening_schedule')) then
161 call json_get(json, '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(json, 'update.enabled', &
169 update_enabled, .false.)
170
171 ! Control the eigenvalue re-estimation; geometry always refreshes.
172 call json_get_or_default(json, 'update.eigs.enabled', &
173 this%refresh_eigs, .true.)
174 call json_get_or_default(json, 'update.eigs.frequency', &
175 this%refresh_eigs_frequency, 20)
176 call json_get_or_default(json, 'update.eigs.warm_start', &
177 this%eigs_warm_start, .true.)
178 call json_get_or_default(json, '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_from_components(&
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_from_components(&
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. The coarse levels compress their
336 ! geometric factors when the fine level does, and keep them in step
337 ! with a moving mesh through recompute_metrics()
338 if (coef%geo_compression) then
339 do i = 1, this%nlvls - 1
340 call this%phmg_hrchy%lvl(i)%coef%enable_geo_compression( &
341 coef%geo_compression_tol)
342 end do
343 call ax_helm_allocator(this%ax, type_name = "standard_compr")
344 else
345 call ax_helm_allocator(this%ax, type_name = "standard")
346 end if
347
348 ! Interpolator Fine + mg levels
349 allocate(this%intrp(this%nlvls - 1))
350 do i = 1, this%nlvls -1
351 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
352 this%phmg_hrchy%lvl(i)%Xh)
353 end do
354
355 ! Coarse space first. Each level maps from level 0.
356 if (this%update_enabled) then
357 allocate(this%crd_intrp(this%nlvls - 1))
358 do i = 1, this%nlvls - 1
359 call this%crd_intrp(i)%init(this%phmg_hrchy%lvl(i)%Xh, &
360 this%phmg_hrchy%lvl(0)%Xh)
361 end do
362 end if
363
364 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
365 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
366 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
367 this%phmg_hrchy%lvl(this%nlvls -1)%bc_projector, &
368 crs_tamg_itrs, crs_tamg_cheby_degree)
369
370 ! update() only refreshes when `lvl(0)%coef%metrics_version` changes.
371 this%last_metrics_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
372
373 ! Hand the eigenvalue policy to every smoother.
374 if (this%update_enabled) then
375 do i = 0, this%nlvls - 1
376 this%phmg_hrchy%lvl(i)%cheby%warm_start_eigs = this%eigs_warm_start
377 this%phmg_hrchy%lvl(i)%cheby%power_its_refresh = &
378 this%power_its_refresh
379 this%phmg_hrchy%lvl(i)%cheby_device%warm_start_eigs = &
380 this%eigs_warm_start
381 this%phmg_hrchy%lvl(i)%cheby_device%power_its_refresh = &
382 this%power_its_refresh
383 end do
384 call this%amg_solver%set_eig_refresh(this%eigs_warm_start, &
385 this%power_its_refresh)
386
387 if (trim(cheby_acc) .eq. "schwarz") then
388 call neko_warning("PHMG: the Schwarz smoother is not refreshed " // &
389 "when the mesh changes. Its local solves stay at the " // &
390 "initial geometry.")
391 end if
392 end if
393
394 end subroutine phmg_init_from_components
395
396 subroutine phmg_free(this)
397 class(phmg_t), intent(inout) :: this
398 integer :: i
399
400 call this%amg_solver%free()
401
402 if (allocated(this%intrp)) then
403 do i = 1, size(this%intrp)
404 call this%intrp(i)%free()
405 end do
406 deallocate(this%intrp)
407 end if
408
409 if (allocated(this%crd_intrp)) then
410 do i = 1, size(this%crd_intrp)
411 call this%crd_intrp(i)%free()
412 end do
413 deallocate(this%crd_intrp)
414 end if
415
416 if (allocated(this%ax)) then
417 call this%ax%free()
418 deallocate(this%ax)
419 end if
420
421 if (allocated(this%phmg_hrchy%lvl)) then
422 do i = lbound(this%phmg_hrchy%lvl, 1), ubound(this%phmg_hrchy%lvl, 1)
423 call this%phmg_hrchy%lvl(i)%r%free()
424 call this%phmg_hrchy%lvl(i)%w%free()
425 call this%phmg_hrchy%lvl(i)%z%free()
426
427 call this%phmg_hrchy%lvl(i)%cheby%free()
428 call this%phmg_hrchy%lvl(i)%cheby_device%free()
429 call this%phmg_hrchy%lvl(i)%jacobi%free()
430 call this%phmg_hrchy%lvl(i)%device_jacobi%free()
431
432 if (allocated(this%phmg_hrchy%lvl(i)%schwarz%work1)) then
433 call this%phmg_hrchy%lvl(i)%schwarz%free()
434 end if
435
436 call this%phmg_hrchy%lvl(i)%bc_projector%free()
437 call this%phmg_hrchy%lvl(i)%bc%free()
438
439 ! Level 0 borrows Xh, dm_Xh, gs_h and coef from the caller,
440 ! all other levels own them
441 if (this%phmg_hrchy%lvl(i)%lvl .gt. 0) then
442 call this%phmg_hrchy%lvl(i)%coef%free()
443 call this%phmg_hrchy%lvl(i)%gs_h%free()
444 call this%phmg_hrchy%lvl(i)%dm_Xh%free()
445 call this%phmg_hrchy%lvl(i)%Xh%free()
446 deallocate(this%phmg_hrchy%lvl(i)%coef)
447 deallocate(this%phmg_hrchy%lvl(i)%gs_h)
448 deallocate(this%phmg_hrchy%lvl(i)%dm_Xh)
449 deallocate(this%phmg_hrchy%lvl(i)%Xh)
450 end if
451
452 nullify(this%phmg_hrchy%lvl(i)%coef)
453 nullify(this%phmg_hrchy%lvl(i)%gs_h)
454 nullify(this%phmg_hrchy%lvl(i)%dm_Xh)
455 nullify(this%phmg_hrchy%lvl(i)%Xh)
456 end do
457 deallocate(this%phmg_hrchy%lvl)
458 end if
459
460 nullify(this%msh)
461
462 end subroutine phmg_free
463
464 subroutine phmg_solve(this, z, r, n)
465 class(phmg_t), intent(inout) :: this
466 integer, intent(in) :: n
467 real(kind=rp), dimension(n), intent(inout) :: z
468 real(kind=rp), dimension(n), intent(inout) :: r
469 type(c_ptr) :: z_d, r_d
470 type(ksp_monitor_t) :: ksp_results
471 integer :: i
472
473 call profiler_start_region('PHMG_solve', 8)
474 associate( mglvl => this%phmg_hrchy%lvl)
475 if (neko_bcknd_device .eq. 1) then
476 z_d = device_get_ptr(z)
477 r_d = device_get_ptr(r)
478 !We should not work with the input
479 call device_copy(mglvl(0)%r%x_d, r_d, n)
480
481 call device_rzero(mglvl(0)%z%x_d, n)
482 call device_rzero(mglvl(0)%w%x_d, n)
483
484 call this%mg_cycle()
485
486 call device_copy(z_d, mglvl(0)%z%x_d, n)
487 else
488 !OCL NORECURRENCE, NOVREC, NOALIAS
489 !DIR$ CONCURRENT
490 !DIR$ IVDEP
491 !GCC$ ivdep
492 !$omp parallel do
493 do i = 1, n
494 !We should not work with the input
495 mglvl(0)%r%x(i,1,1,1) = r(i)
496
497 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
498 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
499 end do
500 !$omp end parallel do
501
502 call this%mg_cycle()
503
504 !OCL NORECURRENCE, NOVREC, NOALIAS
505 !DIR$ CONCURRENT
506 !DIR$ IVDEP
507 !GCC$ ivdep
508 !$omp parallel do
509 do i = 1, n
510 z(i) = mglvl(0)%z%x(i,1,1,1)
511 end do
512 !$omp end parallel do
513 end if
514 end associate
515 call profiler_end_region('PHMG_solve', 8)
516
517 end subroutine phmg_solve
518
520 subroutine phmg_update(this)
521 class(phmg_t), intent(inout) :: this
522 integer :: fine_version
523 logical :: do_eigs
524
525 if (.not. this%update_enabled) return
526
527 fine_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
528
529 ! Mesh has not changed since the last refresh.
530 if (fine_version .eq. this%last_metrics_version) return
531
533
534 call phmg_update_smoother_acc(this)
535
536 do_eigs = (this%refresh_eigs) .and. &
537 (this%refresh_eigs_frequency .gt. 0) .and. &
538 (mod(this%n_refresh, max(this%refresh_eigs_frequency, 1)) &
539 .eq. 0)
540 if (do_eigs) call phmg_update_smoother_eigs(this)
541
542 this%n_refresh = this%n_refresh + 1
543 this%last_metrics_version = fine_version
544
545 end subroutine phmg_update
546
547
551 class(phmg_t), intent(inout) :: this
552 integer :: i
553
554 if (.not. allocated(this%crd_intrp)) then
555 call neko_error("PHMG: update requested but coordinate " // &
556 "interpolators were not initialized.")
557 end if
558
559 call profiler_start_region('PHMG_update_geometry')
560 associate(mg => this%phmg_hrchy%lvl, nelv => this%msh%nelv)
561 do i = 1, this%nlvls - 1
562 call this%crd_intrp(i)%map(mg(i)%dm_Xh%x%x, mg(0)%dm_Xh%x%x, &
563 nelv, mg(i)%Xh)
564 call this%crd_intrp(i)%map(mg(i)%dm_Xh%y%x, mg(0)%dm_Xh%y%x, &
565 nelv, mg(i)%Xh)
566 call this%crd_intrp(i)%map(mg(i)%dm_Xh%z%x, mg(0)%dm_Xh%z%x, &
567 nelv, mg(i)%Xh)
568
569 call mg(i)%coef%recompute_metrics()
570 end do
571 end associate
572 call profiler_end_region('PHMG_update_geometry')
573
574 end subroutine phmg_update_coarse_geometry
575
576
580 class(phmg_t), intent(inout) :: this
581 integer :: i
582
583 call profiler_start_region('PHMG_update_smoother_acc')
584 associate(mg => this%phmg_hrchy%lvl)
585 select case (trim(this%cheby_acc))
586 case ("jacobi")
587 do i = 0, this%nlvls - 1
588 if (neko_bcknd_device .eq. 1) then
589 call mg(i)%device_jacobi%update()
590 else
591 call mg(i)%jacobi%update()
592 end if
593 end do
594 case ("schwarz")
595 ! Not refreshed. The local solves are built by the FDM from the
596 ! element coordinates, and fdm_t has no update path.
597 case default
598 ! No accelerator to refresh.
599 end select
600 end associate
601 call profiler_end_region('PHMG_update_smoother_acc')
602
603 end subroutine phmg_update_smoother_acc
604
605
608 class(phmg_t), intent(inout) :: this
609 integer :: i
610
611 associate(mg => this%phmg_hrchy%lvl)
612 do i = 0, this%nlvls - 1
613 if (neko_bcknd_device .eq. 1) then
614 mg(i)%cheby_device%recompute_eigs = .true.
615 else
616 mg(i)%cheby%recompute_eigs = .true.
617 end if
618 end do
619 end associate
620
621 call this%amg_solver%invalidate_eigs()
622
623 end subroutine phmg_update_smoother_eigs
624
625
626 subroutine phmg_mg_cycle(this)
627 class(phmg_t), intent(inout) :: this
628 type(ksp_monitor_t) :: ksp_results
629 character(len=2) :: lvl_name
630 integer :: lvl, i
631
632 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
633 msh => this%msh, ax => this%Ax)
634 do lvl = 0, this%nlvls-2
635 write(lvl_name, '(I0)') lvl
636 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
637 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
638 !------------!
639 ! SMOOTH !
640 !------------!
641 if (neko_bcknd_device .eq. 1) then
642 mg(lvl)%cheby_device%zero_initial_guess = .true.
643 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
644 r%x, mg(lvl)%dm_Xh%size(), &
645 mg(lvl)%coef, mg(lvl)%bc_projector, &
646 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
647 else
648 mg(lvl)%cheby%zero_initial_guess = .true.
649 ksp_results = mg(lvl)%cheby%solve(ax, z, &
650 r%x, mg(lvl)%dm_Xh%size(), &
651 mg(lvl)%coef, mg(lvl)%bc_projector, &
652 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
653 end if
654
655 !------------!
656 ! Residual !
657 !------------!
658 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
659 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
660 glb_cmd_event)
661 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
662 call mg(lvl)%bc_projector%apply(w%x, mg(lvl)%dm_Xh%size())
663
664 if (neko_bcknd_device .eq. 1) then
665 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
666 else
667 !OCL NORECURRENCE, NOVREC, NOALIAS
668 !DIR$ CONCURRENT
669 !DIR$ IVDEP
670 !GCC$ ivdep
671 !$omp parallel do
672 do i = 1, mg(lvl)%dm_Xh%size()
673 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
674 end do
675 !$omp end parallel do
676 end if
677
678 !------------!
679 ! Restrict !
680 !------------!
681 if (neko_bcknd_device .eq. 1) then
682 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
683 else
684 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
685 end if
686
687 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
688
689 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
690 gs_op_add, glb_cmd_event)
691 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
692
693 call mg(lvl+1)%bc_projector%apply( &
694 mg(lvl+1)%r%x, &
695 mg(lvl+1)%dm_Xh%size())
696
697 if (neko_bcknd_device .eq. 1) then
698 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
699 else
700 !OCL NORECURRENCE, NOVREC, NOALIAS
701 !DIR$ CONCURRENT
702 !DIR$ IVDEP
703 !GCC$ ivdep
704 !$omp parallel do
705 do i = 1, mg(lvl+1)%dm_Xh%size()
706 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
707 end do
708 !$omp end parallel do
709 end if
710 end associate
711 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
712 end do
713
714 call profiler_start_region( 'PHMG_coarse-solve' )
715 !------------!
716 ! SOLVE !
717 !------------!
718 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
719 mg(this%nlvls-1)%r%x, &
720 mg(this%nlvls-1)%dm_Xh%size())
721 call profiler_end_region( 'PHMG_coarse-solve' )
722
723 do lvl = (this%nlvls-2), 0, -1
724 write(lvl_name, '(I0)') lvl
725 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
726 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
727 !------------!
728 ! Project !
729 !------------!
730 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
731
732 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
733 glb_cmd_event)
734 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
735
736 if (neko_bcknd_device .eq. 1) then
737 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
738 else
739 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
740 end if
741
742 !------------!
743 ! Correct !
744 !------------!
745 if (neko_bcknd_device .eq. 1) then
746 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
747 else
748 !OCL NORECURRENCE, NOVREC, NOALIAS
749 !DIR$ CONCURRENT
750 !DIR$ IVDEP
751 !GCC$ ivdep
752 !$omp parallel do
753 do i = 1, mg(lvl)%dm_Xh%size()
754 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
755 end do
756 !$omp end parallel do
757 end if
758
759 !------------!
760 ! SMOOTH !
761 !------------!
762 if (neko_bcknd_device .eq. 1) then
763 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
764 r%x, mg(lvl)%dm_Xh%size(), &
765 mg(lvl)%coef, mg(lvl)%bc_projector, &
766 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
767 else
768 ksp_results = mg(lvl)%cheby%solve(ax, z, &
769 r%x, mg(lvl)%dm_Xh%size(), &
770 mg(lvl)%coef, mg(lvl)%bc_projector, &
771 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
772 end if
773 end associate
774 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
775 end do
776 end associate
777
778 end subroutine phmg_mg_cycle
779
789 subroutine phmg_jacobi_smoother(z, r, w, mg, msh, Ax, n, lvl)
790 type(phmg_lvl_t) :: mg
791 class(ax_t), intent(inout) :: Ax
792 type(mesh_t), intent(inout) :: msh
793 type(field_t), intent(inout) :: z, r, w
794 integer, intent(in) :: n, lvl
795 integer :: i, j, iblk, ni, niblk
796
797 ni = mg%smoother_itrs
798 if (neko_bcknd_device .eq. 1) then
799 do i = 1, ni
800 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
801 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
802 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
803 call mg%bc_projector%apply(w%x, n)
804 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
805
806 call mg%device_jacobi%solve(w%x, w%x, n)
807
808 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
809 end do
810 else
811 do i = 1, ni
812 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
813 call mg%gs_h%op(w%x, n, gs_op_add)
814 call mg%bc_projector%apply(w%x, n)
815 call add2s1(w%x, r%x, -1.0_rp, n)
816
817 call mg%jacobi%solve(w%x, w%x, n)
818
819 call add2s2(z%x, w%x, 0.6_rp, n)
820 end do
821 end if
822 end subroutine phmg_jacobi_smoother
823
824
825 subroutine phmg_resid_monitor(z, r, w, mg, msh, Ax, lvl, typ)
826 integer :: lvl, typ
827 type(phmg_lvl_t) :: mg
828 class(ax_t), intent(inout) :: Ax
829 type(mesh_t), intent(inout) :: msh
830 type(field_t) :: z, r, w
831 real(kind=rp) :: val
832 character(len=LOG_SIZE) :: log_buf
833 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
834 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
835 call mg%bc_projector%apply(w%x, mg%dm_Xh%size())
836 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
837 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
838 if (typ .eq. 1) then
839 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO - PRE', lvl, val
840 else if (typ .eq. 2) then
841 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO -POST', lvl, val
842 else if (typ .eq. 3) then
843 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO- PRE', lvl, val
844 else if (typ .eq. 4) then
845 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO-POST', lvl, val
846 else if (typ .eq. 5) then
847 write(log_buf, '(A15,I4,F12.6)') 'TAMG - PRE', lvl, val
848 else if (typ .eq. 6) then
849 write(log_buf, '(A15,I4,F12.6)') 'TAMG -POST', lvl, val
850 else
851 write(log_buf, '(A15,I4,F12.6)') 'RESID', lvl, val
852 end if
853 call neko_log%message(log_buf)
854 end subroutine phmg_resid_monitor
855
856 subroutine print_phmg_info(nlvls, smoo_type, phmg)
857 integer, intent(in) :: nlvls
858 integer, intent(in) :: smoo_type
859 type(phmg_hrchy_t) :: phmg
860 integer :: i, clvl
861 character(len=LOG_SIZE) :: log_buf, smoo_name
862
863 call neko_log%section('PHMG')
864
865 if (smoo_type .eq. 1) then
866 write(smoo_name, '(A16)') 'CHEBY-acc JACOBI'
867 else if (smoo_type .eq. 2) then
868 write(smoo_name, '(A17)') 'CHEBY-acc SCHWARZ'
869 else
870 write(smoo_name, '(A5)') 'CHEBY'
871 end if
872
873 write(log_buf, '(A28,I2,A8)') &
874 'Creating PHMG hierarchy with', &
875 nlvls, 'levels.'
876 call neko_log%message(log_buf)
877
878 clvl = nlvls - 1
879 do i = 0, nlvls-1
880 write(log_buf, '(A8,I2,A8,I2)') &
881 '-- level', i, '-- lx:', phmg%lvl(i)%Xh%lx
882 call neko_log%message(log_buf)
883
884 if (i .eq. clvl) then
885 write(log_buf, '(A19,A20)') &
886 'Solve:', 'tAMG'
887 call neko_log%message(log_buf)
888 else
889 write(log_buf, '(A22,A20)') &
890 'Smoother:', &
891 trim(smoo_name)
892 call neko_log%message(log_buf)
893
894 write(log_buf, '(A28,I2)') &
895 'Smoother Iters:', &
896 phmg%lvl(i)%smoother_itrs
897 call neko_log%message(log_buf)
898 end if
899 end do
900
901 call neko_log%end_section()
902
903 end subroutine print_phmg_info
904
905end 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:119
integer, parameter, public coef_operator
Retain only what applying a discrete operator needs: , h1, h2, B and mult.
Definition coef.f90:132
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:92
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:580
subroutine phmg_jacobi_smoother(z, r, w, mg, msh, ax, n, lvl)
Wraps jacobi solve as a residual update relaxation method.
Definition phmg.f90:790
subroutine phmg_solve(this, z, r, n)
Definition phmg.f90:465
subroutine phmg_update(this)
Bring the preconditioner back in sync after the mesh has changed.
Definition phmg.f90:521
subroutine phmg_update_smoother_eigs(this)
Mark every smoother to re-estimate its eigenvalues on the next solve.
Definition phmg.f90:608
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:551
subroutine phmg_free(this)
Definition phmg.f90:397
subroutine phmg_mg_cycle(this)
Definition phmg.f90:627
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
Definition phmg.f90:826
subroutine print_phmg_info(nlvls, smoo_type, phmg)
Definition phmg.f90:857
subroutine phmg_init(this, coef, bclst, json)
Definition phmg.f90:135
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:454
Base type for a matrix-vector product providing .
Definition ax.f90:44
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:136
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:47
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:44
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