Neko 1.99.6
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
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(:)
100 type(mesh_t), pointer :: msh
101 contains
102 procedure, pass(this) :: init => phmg_init
103 procedure, pass(this) :: init_from_components => &
105 procedure, pass(this) :: free => phmg_free
106 procedure, pass(this) :: solve => phmg_solve
107 procedure, pass(this) :: update => phmg_update
108 procedure, private, pass(this) :: mg_cycle => phmg_mg_cycle
109 end type phmg_t
110
111contains
112
113 subroutine phmg_init(this, coef, bclst, phmg_params)
114 class(phmg_t), intent(inout), target :: this
115 type(coef_t), intent(in), target :: coef
116 type(bc_list_t), intent(inout), target :: bclst
117 type(json_file), intent(inout) :: phmg_params
118 integer :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
119 integer :: smoother_itrs
120 character(len=:), allocatable :: cheby_acc
121 integer, allocatable :: pcrs_sched(:)
122
123 call json_get_or_default(phmg_params, 'smoother_iterations', &
124 smoother_itrs, 3)
125
126 call json_get_or_default(phmg_params, 'smoother_cheby_acc', &
127 cheby_acc, "jacobi")
128
129 call json_get_or_default(phmg_params, 'coarse_grid.levels', &
130 crs_tamg_lvls, 3)
131
132 call json_get_or_default(phmg_params, 'coarse_grid.iterations', &
133 crs_tamg_itrs, 1)
134
135 call json_get_or_default(phmg_params, 'coarse_grid.cheby_degree', &
136 crs_tamg_cheby_degree, 4)
137
138 if (phmg_params%valid_path('pcoarsening_schedule')) then
139 call json_get(phmg_params, 'pcoarsening_schedule', pcrs_sched)
140 else
141 allocate(pcrs_sched(2))
142 pcrs_sched(1) = 3
143 pcrs_sched(2) = 1
144 end if
145
146
147 call this%init_from_components(coef, bclst, smoother_itrs, &
148 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree,&
149 pcrs_sched)
150
151 end subroutine phmg_init
152
153 subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, &
154 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, &
155 pcrs_sched)
156 class(phmg_t), intent(inout), target :: this
157 type(coef_t), intent(in), target :: coef
158 type(bc_list_t), intent(inout), target :: bclst
159 integer, intent(in) :: smoother_itrs
160 character(len=:), allocatable :: cheby_acc
161 integer, intent(in) :: crs_tamg_lvls, crs_tamg_itrs
162 integer, intent(in) :: crs_tamg_cheby_degree
163 integer, intent(in), allocatable :: pcrs_sched(:)
164 integer :: lx_crs, lx_mid
165 integer, allocatable :: lx_lvls(:)
166 integer :: n, i, j, st
167 class(bc_t), pointer :: bc_j
168 logical :: use_jacobi, use_cheby
169 use_jacobi = .true.
170 use_cheby = .true.
171
172 this%msh => coef%msh
173
174 this%nlvls = size(pcrs_sched) + 1
175 allocate(lx_lvls(0:this%nlvls - 1))
176 lx_lvls(1:) = pcrs_sched + 1
177
178 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
179
180 this%phmg_hrchy%lvl(0)%lvl = 0
181 this%phmg_hrchy%lvl(0)%smoother_itrs = smoother_itrs
182 this%phmg_hrchy%lvl(0)%Xh => coef%Xh
183 this%phmg_hrchy%lvl(0)%coef => coef
184 this%phmg_hrchy%lvl(0)%dm_Xh => coef%dof
185 this%phmg_hrchy%lvl(0)%gs_h => coef%gs_h
186
187 do i = 1, this%nlvls - 1
188 allocate(this%phmg_hrchy%lvl(i)%Xh)
189 allocate(this%phmg_hrchy%lvl(i)%dm_Xh)
190 allocate(this%phmg_hrchy%lvl(i)%gs_h)
191 allocate(this%phmg_hrchy%lvl(i)%coef)
192
193 this%phmg_hrchy%lvl(i)%lvl = i
194 this%phmg_hrchy%lvl(i)%smoother_itrs = smoother_itrs
195 call this%phmg_hrchy%lvl(i)%Xh%init(gll, lx_lvls(i), lx_lvls(i), &
196 lx_lvls(i))
197 call this%phmg_hrchy%lvl(i)%dm_Xh%init(coef%msh, &
198 this%phmg_hrchy%lvl(i)%Xh)
199 call this%phmg_hrchy%lvl(i)%gs_h%init(this%phmg_hrchy%lvl(i)%dm_Xh)
200 call this%phmg_hrchy%lvl(i)%coef%init(this%phmg_hrchy%lvl(i)%gs_h)
201 end do
202
203 do i = 0, this%nlvls - 1
204 call this%phmg_hrchy%lvl(i)%r%init(this%phmg_hrchy%lvl(i)%dm_Xh)
205 call this%phmg_hrchy%lvl(i)%w%init(this%phmg_hrchy%lvl(i)%dm_Xh)
206 call this%phmg_hrchy%lvl(i)%z%init(this%phmg_hrchy%lvl(i)%dm_Xh)
207
208 this%phmg_hrchy%lvl(i)%coef%ifh2 = coef%ifh2
209 call copy(this%phmg_hrchy%lvl(i)%coef%h1, coef%h1, &
210 this%phmg_hrchy%lvl(i)%dm_Xh%size())
211
212 call this%phmg_hrchy%lvl(i)%bc%init_base(this%phmg_hrchy%lvl(i)%coef)
213 if (bclst%size() .gt. 0 ) then
214 do j = 1, bclst%size()
215 bc_j => bclst%get(j)
216 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
217 end do
218 end if
219 call this%phmg_hrchy%lvl(i)%bc%finalize()
220 call this%phmg_hrchy%lvl(i)%bc%set_g(0.0_rp)
221 call this%phmg_hrchy%lvl(i)%bclst%init()
222 call this%phmg_hrchy%lvl(i)%bclst%append(this%phmg_hrchy%lvl(i)%bc)
223
225 if (trim(cheby_acc) .eq. "schwarz") then
226 call this%phmg_hrchy%lvl(i)%schwarz%init( &
227 this%phmg_hrchy%lvl(i)%Xh, &
228 this%phmg_hrchy%lvl(i)%dm_Xh, &
229 this%phmg_hrchy%lvl(i)%gs_h, &
230 this%phmg_hrchy%lvl(i)%bclst, &
231 coef%msh)
232 end if
233
234 if (neko_bcknd_device .eq. 1) then
235 call this%phmg_hrchy%lvl(i)%device_jacobi%init(&
236 this%phmg_hrchy%lvl(i)%coef, &
237 this%phmg_hrchy%lvl(i)%dm_Xh, &
238 this%phmg_hrchy%lvl(i)%gs_h)
239 else
240 call this%phmg_hrchy%lvl(i)%jacobi%init(&
241 this%phmg_hrchy%lvl(i)%coef, &
242 this%phmg_hrchy%lvl(i)%dm_Xh, &
243 this%phmg_hrchy%lvl(i)%gs_h)
244 end if
245
246 if (neko_bcknd_device .eq. 1) then
247 if (trim(cheby_acc) .eq. "jacobi") then
248 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
249 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
250 this%phmg_hrchy%lvl(i)%device_jacobi)
251 st = 1
252 else
253 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
254 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
255 st = 0
256 if (trim(cheby_acc) .eq. "schwarz") then
257 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
258 this%phmg_hrchy%lvl(i)%schwarz
259 st = 2
260 end if
261 end if
262 else
263 if (trim(cheby_acc) .eq. "jacobi") then
264 call this%phmg_hrchy%lvl(i)%cheby%init( &
265 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
266 this%phmg_hrchy%lvl(i)%jacobi)
267 st = 1
268 else
269 call this%phmg_hrchy%lvl(i)%cheby%init( &
270 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
271 st = 0
272 if (trim(cheby_acc) .eq. "schwarz") then
273 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
274 this%phmg_hrchy%lvl(i)%schwarz
275 st = 2
276 end if
277 end if
278 end if
279
280 end do
281
282 call print_phmg_info(this%nlvls, st, this%phmg_hrchy)
283
284 ! Create backend specific Ax operator
285 call ax_helm_factory(this%ax, full_formulation = .false.)
286
287 ! Interpolator Fine + mg levels
288 allocate(this%intrp(this%nlvls - 1))
289 do i = 1, this%nlvls -1
290 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
291 this%phmg_hrchy%lvl(i)%Xh)
292 end do
293
294 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
295 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
296 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
297 this%phmg_hrchy%lvl(this%nlvls -1)%bclst, &
298 crs_tamg_itrs, crs_tamg_cheby_degree)
299
300 end subroutine phmg_init_from_components
301
302 subroutine phmg_free(this)
303 class(phmg_t), intent(inout) :: this
304 integer :: i
305
306 call this%amg_solver%free()
307
308 if (allocated(this%intrp)) then
309 do i = 1, size(this%intrp)
310 call this%intrp(i)%free()
311 end do
312 deallocate(this%intrp)
313 end if
314
315 if (allocated(this%ax)) then
316 deallocate(this%ax)
317 end if
318
319 if (allocated(this%phmg_hrchy%lvl)) then
320 do i = lbound(this%phmg_hrchy%lvl, 1), ubound(this%phmg_hrchy%lvl, 1)
321 call this%phmg_hrchy%lvl(i)%r%free()
322 call this%phmg_hrchy%lvl(i)%w%free()
323 call this%phmg_hrchy%lvl(i)%z%free()
324
325 call this%phmg_hrchy%lvl(i)%cheby%free()
326 call this%phmg_hrchy%lvl(i)%cheby_device%free()
327 call this%phmg_hrchy%lvl(i)%jacobi%free()
328 call this%phmg_hrchy%lvl(i)%device_jacobi%free()
329
330 if (allocated(this%phmg_hrchy%lvl(i)%schwarz%work1)) then
331 call this%phmg_hrchy%lvl(i)%schwarz%free()
332 end if
333
334 call this%phmg_hrchy%lvl(i)%bclst%free()
335 call this%phmg_hrchy%lvl(i)%bc%free()
336
337 ! Level 0 borrows Xh, dm_Xh, gs_h and coef from the caller,
338 ! all other levels own them
339 if (this%phmg_hrchy%lvl(i)%lvl .gt. 0) then
340 call this%phmg_hrchy%lvl(i)%coef%free()
341 call this%phmg_hrchy%lvl(i)%gs_h%free()
342 call this%phmg_hrchy%lvl(i)%dm_Xh%free()
343 call this%phmg_hrchy%lvl(i)%Xh%free()
344 deallocate(this%phmg_hrchy%lvl(i)%coef)
345 deallocate(this%phmg_hrchy%lvl(i)%gs_h)
346 deallocate(this%phmg_hrchy%lvl(i)%dm_Xh)
347 deallocate(this%phmg_hrchy%lvl(i)%Xh)
348 end if
349
350 nullify(this%phmg_hrchy%lvl(i)%coef)
351 nullify(this%phmg_hrchy%lvl(i)%gs_h)
352 nullify(this%phmg_hrchy%lvl(i)%dm_Xh)
353 nullify(this%phmg_hrchy%lvl(i)%Xh)
354 end do
355 deallocate(this%phmg_hrchy%lvl)
356 end if
357
358 nullify(this%msh)
359
360 end subroutine phmg_free
361
362 subroutine phmg_solve(this, z, r, n)
363 class(phmg_t), intent(inout) :: this
364 integer, intent(in) :: n
365 real(kind=rp), dimension(n), intent(inout) :: z
366 real(kind=rp), dimension(n), intent(inout) :: r
367 type(c_ptr) :: z_d, r_d
368 type(ksp_monitor_t) :: ksp_results
369 integer :: i
370
371 call profiler_start_region('PHMG_solve', 8)
372 associate( mglvl => this%phmg_hrchy%lvl)
373 if (neko_bcknd_device .eq. 1) then
374 z_d = device_get_ptr(z)
375 r_d = device_get_ptr(r)
376 !We should not work with the input
377 call device_copy(mglvl(0)%r%x_d, r_d, n)
378
379 call device_rzero(mglvl(0)%z%x_d, n)
380 call device_rzero(mglvl(0)%w%x_d, n)
381
382 call this%mg_cycle()
383
384 call device_copy(z_d, mglvl(0)%z%x_d, n)
385 else
386 !OCL NORECURRENCE, NOVREC, NOALIAS
387 !DIR$ CONCURRENT
388 !DIR$ IVDEP
389 !GCC$ ivdep
390 !$omp parallel do
391 do i = 1, n
392 !We should not work with the input
393 mglvl(0)%r%x(i,1,1,1) = r(i)
394
395 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
396 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
397 end do
398 !$omp end parallel do
399
400 call this%mg_cycle()
401
402 !OCL NORECURRENCE, NOVREC, NOALIAS
403 !DIR$ CONCURRENT
404 !DIR$ IVDEP
405 !GCC$ ivdep
406 !$omp parallel do
407 do i = 1, n
408 z(i) = mglvl(0)%z%x(i,1,1,1)
409 end do
410 !$omp end parallel do
411 end if
412 end associate
413 call profiler_end_region('PHMG_solve', 8)
414
415 end subroutine phmg_solve
416
417 subroutine phmg_update(this)
418 class(phmg_t), intent(inout) :: this
419 end subroutine phmg_update
420
421
422 subroutine phmg_mg_cycle(this)
423 class(phmg_t), intent(inout) :: this
424 type(ksp_monitor_t) :: ksp_results
425 character(len=2) :: lvl_name
426 integer :: lvl, i
427
428 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
429 msh => this%msh, ax => this%Ax)
430 do lvl = 0, this%nlvls-2
431 write(lvl_name, '(I0)') lvl
432 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
433 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
434 !------------!
435 ! SMOOTH !
436 !------------!
437 if (neko_bcknd_device .eq. 1) then
438 mg(lvl)%cheby_device%zero_initial_guess = .true.
439 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
440 r%x, mg(lvl)%dm_Xh%size(), &
441 mg(lvl)%coef, mg(lvl)%bclst, &
442 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
443 else
444 mg(lvl)%cheby%zero_initial_guess = .true.
445 ksp_results = mg(lvl)%cheby%solve(ax, z, &
446 r%x, mg(lvl)%dm_Xh%size(), &
447 mg(lvl)%coef, mg(lvl)%bclst, &
448 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
449 end if
450
451 !------------!
452 ! Residual !
453 !------------!
454 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
455 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
456 glb_cmd_event)
457 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
458 call mg(lvl)%bclst%apply_scalar(w%x, mg(lvl)%dm_Xh%size())
459
460 if (neko_bcknd_device .eq. 1) then
461 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
462 else
463 !OCL NORECURRENCE, NOVREC, NOALIAS
464 !DIR$ CONCURRENT
465 !DIR$ IVDEP
466 !GCC$ ivdep
467 !$omp parallel do
468 do i = 1, mg(lvl)%dm_Xh%size()
469 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
470 end do
471 !$omp end parallel do
472 end if
473
474 !------------!
475 ! Restrict !
476 !------------!
477 if (neko_bcknd_device .eq. 1) then
478 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
479 else
480 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
481 end if
482
483 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
484
485 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
486 gs_op_add, glb_cmd_event)
487 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
488
489 call mg(lvl+1)%bclst%apply_scalar( &
490 mg(lvl+1)%r%x, &
491 mg(lvl+1)%dm_Xh%size())
492
493 if (neko_bcknd_device .eq. 1) then
494 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
495 else
496 !OCL NORECURRENCE, NOVREC, NOALIAS
497 !DIR$ CONCURRENT
498 !DIR$ IVDEP
499 !GCC$ ivdep
500 !$omp parallel do
501 do i = 1, mg(lvl+1)%dm_Xh%size()
502 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
503 end do
504 !$omp end parallel do
505 end if
506 end associate
507 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
508 end do
509
510 call profiler_start_region( 'PHMG_coarse-solve' )
511 !------------!
512 ! SOLVE !
513 !------------!
514 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
515 mg(this%nlvls-1)%r%x, &
516 mg(this%nlvls-1)%dm_Xh%size())
517 call profiler_end_region( 'PHMG_coarse-solve' )
518
519 do lvl = (this%nlvls-2), 0, -1
520 write(lvl_name, '(I0)') lvl
521 call profiler_start_region( "PHMG_level_" // trim(lvl_name))
522 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
523 !------------!
524 ! Project !
525 !------------!
526 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
527
528 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
529 glb_cmd_event)
530 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
531
532 if (neko_bcknd_device .eq. 1) then
533 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
534 else
535 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
536 end if
537
538 !------------!
539 ! Correct !
540 !------------!
541 if (neko_bcknd_device .eq. 1) then
542 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
543 else
544 !OCL NORECURRENCE, NOVREC, NOALIAS
545 !DIR$ CONCURRENT
546 !DIR$ IVDEP
547 !GCC$ ivdep
548 !$omp parallel do
549 do i = 1, mg(lvl)%dm_Xh%size()
550 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
551 end do
552 !$omp end parallel do
553 end if
554
555 !------------!
556 ! SMOOTH !
557 !------------!
558 if (neko_bcknd_device .eq. 1) then
559 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
560 r%x, mg(lvl)%dm_Xh%size(), &
561 mg(lvl)%coef, mg(lvl)%bclst, &
562 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
563 else
564 ksp_results = mg(lvl)%cheby%solve(ax, z, &
565 r%x, mg(lvl)%dm_Xh%size(), &
566 mg(lvl)%coef, mg(lvl)%bclst, &
567 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
568 end if
569 end associate
570 call profiler_end_region( "PHMG_level_" // trim(lvl_name))
571 end do
572 end associate
573
574 end subroutine phmg_mg_cycle
575
585 subroutine phmg_jacobi_smoother(z, r, w, mg, msh, Ax, n, lvl)
586 type(phmg_lvl_t) :: mg
587 class(ax_t), intent(inout) :: Ax
588 type(mesh_t), intent(inout) :: msh
589 type(field_t), intent(inout) :: z, r, w
590 integer, intent(in) :: n, lvl
591 integer :: i, j, iblk, ni, niblk
592
593 ni = mg%smoother_itrs
594 if (neko_bcknd_device .eq. 1) then
595 do i = 1, ni
596 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
597 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
598 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
599 call mg%bclst%apply_scalar(w%x, n)
600 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
601
602 call mg%device_jacobi%solve(w%x, w%x, n)
603
604 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
605 end do
606 else
607 do i = 1, ni
608 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
609 call mg%gs_h%op(w%x, n, gs_op_add)
610 call mg%bclst%apply_scalar(w%x, n)
611 call add2s1(w%x, r%x, -1.0_rp, n)
612
613 call mg%jacobi%solve(w%x, w%x, n)
614
615 call add2s2(z%x, w%x, 0.6_rp, n)
616 end do
617 end if
618 end subroutine phmg_jacobi_smoother
619
620
621 subroutine phmg_resid_monitor(z, r, w, mg, msh, Ax, lvl, typ)
622 integer :: lvl, typ
623 type(phmg_lvl_t) :: mg
624 class(ax_t), intent(inout) :: Ax
625 type(mesh_t), intent(inout) :: msh
626 type(field_t) :: z, r, w
627 real(kind=rp) :: val
628 character(len=LOG_SIZE) :: log_buf
629 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
630 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
631 call mg%bclst%apply_scalar(w%x, mg%dm_Xh%size())
632 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
633 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
634 if (typ .eq. 1) then
635 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO - PRE', lvl, val
636 else if (typ .eq. 2) then
637 write(log_buf, '(A15,I4,F12.6)') 'PRESMOO -POST', lvl, val
638 else if (typ .eq. 3) then
639 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO- PRE', lvl, val
640 else if (typ .eq. 4) then
641 write(log_buf, '(A15,I4,F12.6)') 'POSTSMOO-POST', lvl, val
642 else if (typ .eq. 5) then
643 write(log_buf, '(A15,I4,F12.6)') 'TAMG - PRE', lvl, val
644 else if (typ .eq. 6) then
645 write(log_buf, '(A15,I4,F12.6)') 'TAMG -POST', lvl, val
646 else
647 write(log_buf, '(A15,I4,F12.6)') 'RESID', lvl, val
648 end if
649 call neko_log%message(log_buf)
650 end subroutine phmg_resid_monitor
651
652 subroutine print_phmg_info(nlvls, smoo_type, phmg)
653 integer, intent(in) :: nlvls
654 integer, intent(in) :: smoo_type
655 type(phmg_hrchy_t) :: phmg
656 integer :: i, clvl
657 character(len=LOG_SIZE) :: log_buf, smoo_name
658
659 call neko_log%section('PHMG')
660
661 if (smoo_type .eq. 1) then
662 write(smoo_name, '(A16)') 'CHEBY-acc JACOBI'
663 else if (smoo_type .eq. 2) then
664 write(smoo_name, '(A17)') 'CHEBY-acc SCHWARZ'
665 else
666 write(smoo_name, '(A5)') 'CHEBY'
667 end if
668
669 write(log_buf, '(A28,I2,A8)') &
670 'Creating PHMG hierarchy with', &
671 nlvls, 'levels.'
672 call neko_log%message(log_buf)
673
674 clvl = nlvls - 1
675 do i = 0, nlvls-1
676 write(log_buf, '(A8,I2,A8,I2)') &
677 '-- level', i, '-- lx:', phmg%lvl(i)%Xh%lx
678 call neko_log%message(log_buf)
679
680 if (i .eq. clvl) then
681 write(log_buf, '(A19,A20)') &
682 'Solve:', 'tAMG'
683 call neko_log%message(log_buf)
684 else
685 write(log_buf, '(A22,A20)') &
686 'Smoother:', &
687 trim(smoo_name)
688 call neko_log%message(log_buf)
689
690 write(log_buf, '(A28,I2)') &
691 'Smoother Iters:', &
692 phmg%lvl(i)%smoother_itrs
693 call neko_log%message(log_buf)
694 end if
695 end do
696
697 call neko_log%end_section()
698
699 end subroutine print_phmg_info
700
701end 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_jacobi_smoother(z, r, w, mg, msh, ax, n, lvl)
Wraps jacobi solve as a residual update relaxation method.
Definition phmg.f90:586
subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, pcrs_sched)
Definition phmg.f90:156
subroutine phmg_solve(this, z, r, n)
Definition phmg.f90:363
subroutine phmg_init(this, coef, bclst, phmg_params)
Definition phmg.f90:114
subroutine phmg_update(this)
Definition phmg.f90:418
subroutine phmg_free(this)
Definition phmg.f90:303
subroutine phmg_mg_cycle(this)
Definition phmg.f90:423
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
Definition phmg.f90:622
subroutine print_phmg_info(nlvls, smoo_type, phmg)
Definition phmg.f90:653
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
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:52
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.