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