55 use json_module,
only : json_file
67 use,
intrinsic :: iso_c_binding
74 integer :: smoother_itrs = 10
77 type(
gs_t),
pointer :: gs_h
98 class(
ax_t),
allocatable :: ax
103 procedure, pass(this) :: init_from_components => &
113 class(
phmg_t),
intent(inout),
target :: this
114 type(
coef_t),
intent(in),
target :: coef
115 type(
bc_list_t),
intent(inout),
target :: bclst
116 type(json_file),
intent(inout) :: phmg_params
117 integer :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
118 integer :: smoother_itrs
119 character(len=:),
allocatable :: cheby_acc
120 integer,
allocatable :: pcrs_sched(:)
135 crs_tamg_cheby_degree, 5)
137 if (phmg_params%valid_path(
'pcoarsening_schedule'))
then
138 call json_get(phmg_params,
'pcoarsening_schedule', pcrs_sched)
140 allocate(pcrs_sched(2))
146 call this%init_from_components(coef, bclst, smoother_itrs, &
147 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree,&
153 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, &
155 class(
phmg_t),
intent(inout),
target :: this
156 type(
coef_t),
intent(in),
target :: coef
157 type(
bc_list_t),
intent(inout),
target :: bclst
158 integer,
intent(in) :: smoother_itrs
159 character(len=:),
allocatable :: cheby_acc
160 integer,
intent(in) :: crs_tamg_lvls, crs_tamg_itrs
161 integer,
intent(in) :: crs_tamg_cheby_degree
162 integer,
intent(in),
allocatable :: pcrs_sched(:)
163 integer :: lx_crs, lx_mid
164 integer,
allocatable :: lx_lvls(:)
165 integer :: n, i, j, st
166 class(
bc_t),
pointer :: bc_j
167 logical :: use_jacobi, use_cheby
173 this%nlvls =
size(pcrs_sched) + 1
174 allocate(lx_lvls(0:this%nlvls - 1))
175 lx_lvls(1:) = pcrs_sched + 1
177 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
179 this%phmg_hrchy%lvl(0)%lvl = 0
180 this%phmg_hrchy%lvl(0)%smoother_itrs = smoother_itrs
181 this%phmg_hrchy%lvl(0)%Xh => coef%Xh
182 this%phmg_hrchy%lvl(0)%coef => coef
183 this%phmg_hrchy%lvl(0)%dm_Xh => coef%dof
184 this%phmg_hrchy%lvl(0)%gs_h => coef%gs_h
186 do i = 1, this%nlvls - 1
187 allocate(this%phmg_hrchy%lvl(i)%Xh)
188 allocate(this%phmg_hrchy%lvl(i)%dm_Xh)
189 allocate(this%phmg_hrchy%lvl(i)%gs_h)
190 allocate(this%phmg_hrchy%lvl(i)%coef)
192 this%phmg_hrchy%lvl(i)%lvl = i
193 this%phmg_hrchy%lvl(i)%smoother_itrs = smoother_itrs
194 call this%phmg_hrchy%lvl(i)%Xh%init(
gll, lx_lvls(i), lx_lvls(i), &
196 call this%phmg_hrchy%lvl(i)%dm_Xh%init(coef%msh, &
197 this%phmg_hrchy%lvl(i)%Xh)
198 call this%phmg_hrchy%lvl(i)%gs_h%init(this%phmg_hrchy%lvl(i)%dm_Xh)
199 call this%phmg_hrchy%lvl(i)%coef%init(this%phmg_hrchy%lvl(i)%gs_h)
202 do i = 0, this%nlvls - 1
203 call this%phmg_hrchy%lvl(i)%r%init(this%phmg_hrchy%lvl(i)%dm_Xh)
204 call this%phmg_hrchy%lvl(i)%w%init(this%phmg_hrchy%lvl(i)%dm_Xh)
205 call this%phmg_hrchy%lvl(i)%z%init(this%phmg_hrchy%lvl(i)%dm_Xh)
207 this%phmg_hrchy%lvl(i)%coef%ifh2 = coef%ifh2
208 call copy(this%phmg_hrchy%lvl(i)%coef%h1, coef%h1, &
209 this%phmg_hrchy%lvl(i)%dm_Xh%size())
211 call this%phmg_hrchy%lvl(i)%bc%init_base(this%phmg_hrchy%lvl(i)%coef)
212 if (bclst%size() .gt. 0 )
then
213 do j = 1, bclst%size()
215 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
218 call this%phmg_hrchy%lvl(i)%bc%finalize()
219 call this%phmg_hrchy%lvl(i)%bc%set_g(0.0_rp)
220 call this%phmg_hrchy%lvl(i)%bclst%init()
221 call this%phmg_hrchy%lvl(i)%bclst%append(this%phmg_hrchy%lvl(i)%bc)
224 if (trim(cheby_acc) .eq.
"schwarz")
then
225 call this%phmg_hrchy%lvl(i)%schwarz%init( &
226 this%phmg_hrchy%lvl(i)%Xh, &
227 this%phmg_hrchy%lvl(i)%dm_Xh, &
228 this%phmg_hrchy%lvl(i)%gs_h, &
229 this%phmg_hrchy%lvl(i)%bclst, &
234 call this%phmg_hrchy%lvl(i)%device_jacobi%init(&
235 this%phmg_hrchy%lvl(i)%coef, &
236 this%phmg_hrchy%lvl(i)%dm_Xh, &
237 this%phmg_hrchy%lvl(i)%gs_h)
239 call this%phmg_hrchy%lvl(i)%jacobi%init(&
240 this%phmg_hrchy%lvl(i)%coef, &
241 this%phmg_hrchy%lvl(i)%dm_Xh, &
242 this%phmg_hrchy%lvl(i)%gs_h)
246 if (trim(cheby_acc) .eq.
"jacobi")
then
247 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
248 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
249 this%phmg_hrchy%lvl(i)%device_jacobi)
252 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
253 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
255 if (trim(cheby_acc) .eq.
"schwarz")
then
256 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
257 this%phmg_hrchy%lvl(i)%schwarz
262 if (trim(cheby_acc) .eq.
"jacobi")
then
263 call this%phmg_hrchy%lvl(i)%cheby%init( &
264 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
265 this%phmg_hrchy%lvl(i)%jacobi)
268 call this%phmg_hrchy%lvl(i)%cheby%init( &
269 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
271 if (trim(cheby_acc) .eq.
"schwarz")
then
272 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
273 this%phmg_hrchy%lvl(i)%schwarz
284 call ax_helm_factory(this%ax, full_formulation = .false.)
287 allocate(this%intrp(this%nlvls - 1))
288 do i = 1, this%nlvls -1
289 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
290 this%phmg_hrchy%lvl(i)%Xh)
293 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
294 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
295 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
296 this%phmg_hrchy%lvl(this%nlvls -1)%bclst, &
297 crs_tamg_itrs, crs_tamg_cheby_degree)
302 class(
phmg_t),
intent(inout) :: this
306 class(
phmg_t),
intent(inout) :: this
307 integer,
intent(in) :: n
308 real(kind=
rp),
dimension(n),
intent(inout) :: z
309 real(kind=
rp),
dimension(n),
intent(inout) :: r
310 type(c_ptr) :: z_d, r_d
314 associate( mglvl => this%phmg_hrchy%lvl)
323 this%nlvls -1, mglvl, this%intrp, this%msh, this%Ax, &
329 call copy(mglvl(0)%r%x, r, n)
331 mglvl(0)%z%x = 0.0_rp
332 mglvl(0)%w%x = 0.0_rp
335 this%nlvls -1, mglvl, this%intrp, this%msh, this%Ax, &
338 call copy(z, mglvl(0)%z%x, n)
345 class(
phmg_t),
intent(inout) :: this
350 mg, intrp, msh, Ax, amg_solver)
351 type(ksp_monitor_t) :: ksp_results
354 type(interpolator_t) :: intrp(1:clvl)
355 type(tamg_solver_t),
intent(inout) :: amg_solver
356 class(ax_t),
intent(inout) :: ax
357 type(mesh_t),
intent(inout) :: msh
358 type(field_t) :: z, r, w
360 logical :: use_jacobi
364 call profiler_start_region(
'PHMG_cycle', 8)
368 call profiler_start_region(
'PHMG_PreSmooth', 9)
371 mg(lvl)%dm_Xh%size(), lvl)
373 if (neko_bcknd_device .eq. 1)
then
374 mg(lvl)%cheby_device%zero_initial_guess = .true.
375 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
376 r%x, mg(lvl)%dm_Xh%size(), &
377 mg(lvl)%coef, mg(lvl)%bclst, &
378 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
380 mg(lvl)%cheby%zero_initial_guess = .true.
381 ksp_results = mg(lvl)%cheby%solve(ax, z, &
382 r%x, mg(lvl)%dm_Xh%size(), &
383 mg(lvl)%coef, mg(lvl)%bclst, &
384 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
387 call profiler_end_region(
'PHMG_PreSmooth', 9)
392 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
393 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, glb_cmd_event)
394 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
395 call mg(lvl)%bclst%apply_scalar(w%x, mg(lvl)%dm_Xh%size())
397 if (neko_bcknd_device .eq. 1)
then
398 call device_sub3(w%x_d, r%x_d, w%x_d, mg(lvl)%dm_Xh%size())
406 if (neko_bcknd_device .eq. 1)
then
407 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
409 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
412 call profiler_start_region(
'PHMG_map_to_coarse', 9)
413 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
414 call profiler_end_region(
'PHMG_map_to_coarse', 9)
416 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
417 gs_op_add, glb_cmd_event)
418 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
420 call mg(lvl+1)%bclst%apply_scalar( &
422 mg(lvl+1)%dm_Xh%size())
426 if (neko_bcknd_device .eq. 1)
then
427 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
429 mg(lvl+1)%z%x = 0.0_rp
431 if (lvl+1 .eq. clvl)
then
432 call profiler_start_region(
'PHMG_tAMG_coarse_grid', 9)
433 call amg_solver%solve(mg(lvl+1)%z%x, &
435 mg(lvl+1)%dm_Xh%size())
436 call profiler_end_region(
'PHMG_tAMG_coarse_grid', 9)
439 call phmg_mg_cycle(mg(lvl+1)%z, mg(lvl+1)%r, mg(lvl+1)%w, lvl+1, &
440 clvl, mg, intrp, msh, ax, amg_solver)
446 call profiler_start_region(
'PHMG_map_to_fine', 9)
447 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
448 call profiler_end_region(
'PHMG_map_to_fine', 9)
450 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, glb_cmd_event)
451 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
453 if (neko_bcknd_device .eq. 1)
then
454 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
456 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
462 if (neko_bcknd_device .eq. 1)
then
463 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
471 call profiler_start_region(
'PHMG_PostSmooth', 9)
474 mg(lvl)%dm_Xh%size(), lvl)
476 if (neko_bcknd_device .eq. 1)
then
477 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
478 r%x, mg(lvl)%dm_Xh%size(), &
479 mg(lvl)%coef, mg(lvl)%bclst, &
480 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
482 ksp_results = mg(lvl)%cheby%solve(ax, z, &
483 r%x, mg(lvl)%dm_Xh%size(), &
484 mg(lvl)%coef, mg(lvl)%bclst, &
485 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
488 call profiler_end_region(
'PHMG_PostSmooth', 9)
490 call profiler_end_region(
'PHMG_cycle', 8)
504 class(ax_t),
intent(inout) :: Ax
505 type(mesh_t),
intent(inout) :: msh
506 type(field_t),
intent(inout) :: z, r, w
507 integer,
intent(in) :: n, lvl
508 integer :: i, iblk, ni, niblk
510 ni = mg%smoother_itrs
511 if (neko_bcknd_device .eq. 1)
then
513 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
514 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
515 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
516 call mg%bclst%apply_scalar(w%x, n)
517 call device_sub3(w%x_d, r%x_d, w%x_d, n)
519 call mg%device_jacobi%solve(w%x, w%x, n)
521 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
525 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
526 call mg%gs_h%op(w%x, n, gs_op_add)
527 call mg%bclst%apply_scalar(w%x, n)
528 call sub3(w%x, r%x, w%x, n)
530 call mg%jacobi%solve(w%x, w%x, n)
532 call add2s2(z%x, w%x, 0.6_rp, n)
541 class(ax_t),
intent(inout) :: Ax
542 type(mesh_t),
intent(inout) :: msh
543 type(field_t) :: z, r, w
545 character(len=LOG_SIZE) :: log_buf
546 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
547 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
548 call mg%bclst%apply_scalar(w%x, mg%dm_Xh%size())
549 call device_sub3(w%x_d, r%x_d, w%x_d, mg%dm_Xh%size())
550 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
552 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO - PRE', lvl, val
553 else if (typ .eq. 2)
then
554 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO -POST', lvl, val
555 else if (typ .eq. 3)
then
556 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO- PRE', lvl, val
557 else if (typ .eq. 4)
then
558 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO-POST', lvl, val
559 else if (typ .eq. 5)
then
560 write(log_buf,
'(A15,I4,F12.6)')
'TAMG - PRE', lvl, val
561 else if (typ .eq. 6)
then
562 write(log_buf,
'(A15,I4,F12.6)')
'TAMG -POST', lvl, val
564 write(log_buf,
'(A15,I4,F12.6)')
'RESID', lvl, val
566 call neko_log%message(log_buf)
570 integer,
intent(in) :: nlvls
571 integer,
intent(in) :: smoo_type
574 character(len=LOG_SIZE) :: log_buf, smoo_name
576 call neko_log%section(
'PHMG')
578 if (smoo_type .eq. 1)
then
579 write(smoo_name,
'(A16)')
'CHEBY-acc JACOBI'
580 else if (smoo_type .eq. 2)
then
581 write(smoo_name,
'(A17)')
'CHEBY-acc SCHWARZ'
583 write(smoo_name,
'(A5)')
'CHEBY'
586 write(log_buf,
'(A28,I2,A8)') &
587 'Creating PHMG hierarchy with', &
589 call neko_log%message(log_buf)
593 write(log_buf,
'(A8,I2,A8,I2)') &
594 '-- level', i,
'-- lx:',
phmg%lvl(i)%Xh%lx
595 call neko_log%message(log_buf)
597 if (i .eq. clvl)
then
598 write(log_buf,
'(A19,A20)') &
600 call neko_log%message(log_buf)
602 write(log_buf,
'(A22,A20)') &
605 call neko_log%message(log_buf)
607 write(log_buf,
'(A28,I2)') &
609 phmg%lvl(i)%smoother_itrs
610 call neko_log%message(log_buf)
614 call neko_log%end_section()
__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.
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.
Defines a boundary condition.
Chebyshev preconditioner.
Chebyshev preconditioner.
Jacobi preconditioner accelerator backend.
subroutine, public device_sub3(a_d, b_d, c_d, n, strm)
Vector subtraction .
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.
subroutine, public device_stream_wait_event(stream, event, flags)
Synchronize a device stream with an event.
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
type(c_ptr), bind(C), public glb_cmd_event
Event for the global command queue.
Defines a dirichlet boundary condition.
Defines a mapping of the degrees of freedom.
Routines to interpolate between different spaces.
Utilities for retrieving parameters from the case files.
Implements the base abstract type for Krylov solvers plus helper types.
integer, parameter, public ksp_max_iter
Maximum number of iters.
type(log_t), public neko_log
Global log stream.
integer, parameter, public log_size
subroutine, public sub3(a, b, c, n)
Vector subtraction .
subroutine, public add2(a, b, n)
Vector addition .
subroutine, public col2(a, b, n)
Vector multiplication .
subroutine, public copy(a, b, n)
Copy a vector .
subroutine, public add2s2(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on second argument)
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Hybrid ph-multigrid preconditioner.
subroutine phmg_jacobi_smoother(z, r, w, mg, msh, ax, n, lvl)
Wraps jacobi solve as a residual update relaxation method.
subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, pcrs_sched)
subroutine phmg_solve(this, z, r, n)
subroutine phmg_init(this, coef, bclst, phmg_params)
subroutine phmg_update(this)
subroutine phmg_free(this)
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
subroutine print_phmg_info(nlvls, smoo_type, phmg)
recursive subroutine phmg_mg_cycle(z, r, w, lvl, clvl, mg, intrp, msh, ax, amg_solver)
subroutine, public profiler_start_region(name, region_id)
Started a named (name) profiler region.
subroutine, public profiler_end_region(name, region_id)
End the most recently started profiler region.
Overlapping schwarz solves.
Defines a function space.
integer, parameter, public gll
Implements multigrid using the TreeAMG hierarchy structure. USE:
Base type for a matrix-vector product providing .
Base type for a boundary condition.
A list of allocatable `bc_t`. Follows the standard interface of lists.
Defines a Chebyshev preconditioner.
Defines a Chebyshev preconditioner.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Defines a jacobi preconditioner.
Generic Dirichlet boundary condition on .
Interpolation between two space::space_t.
Defines a jacobi preconditioner.
Type for storing initial and final residuals in a Krylov solver.
Base abstract type for a canonical Krylov method, solving .
Defines a canonical Krylov preconditioner.
The function space for the SEM solution fields.
Type for the TreeAMG solver.