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
134 crs_tamg_cheby_degree, 5)
136 call this%init_from_components(coef, bclst, smoother_itrs, &
138 crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
144 crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
145 class(
phmg_t),
intent(inout),
target :: this
146 type(
coef_t),
intent(in),
target :: coef
147 type(
bc_list_t),
intent(inout),
target :: bclst
148 integer,
intent(in) :: smoother_itrs
149 character(len=:),
allocatable :: cheby_acc
150 integer,
intent(in) :: crs_tamg_lvls, crs_tamg_itrs
151 integer,
intent(in) :: crs_tamg_cheby_degree
152 integer :: lx_crs, lx_mid
153 integer,
allocatable :: lx_lvls(:)
154 integer :: n, i, j, st
155 class(
bc_t),
pointer :: bc_j
156 logical :: use_jacobi, use_cheby
166 allocate(lx_lvls(0:this%nlvls - 1))
174 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
176 this%phmg_hrchy%lvl(0)%lvl = 0
177 this%phmg_hrchy%lvl(0)%smoother_itrs = smoother_itrs
178 this%phmg_hrchy%lvl(0)%Xh => coef%Xh
179 this%phmg_hrchy%lvl(0)%coef => coef
180 this%phmg_hrchy%lvl(0)%dm_Xh => coef%dof
181 this%phmg_hrchy%lvl(0)%gs_h => coef%gs_h
183 do i = 1, this%nlvls - 1
184 allocate(this%phmg_hrchy%lvl(i)%Xh)
185 allocate(this%phmg_hrchy%lvl(i)%dm_Xh)
186 allocate(this%phmg_hrchy%lvl(i)%gs_h)
187 allocate(this%phmg_hrchy%lvl(i)%coef)
189 this%phmg_hrchy%lvl(i)%lvl = i
190 this%phmg_hrchy%lvl(i)%smoother_itrs = smoother_itrs
191 call this%phmg_hrchy%lvl(i)%Xh%init(
gll, lx_lvls(i), lx_lvls(i), &
193 call this%phmg_hrchy%lvl(i)%dm_Xh%init(coef%msh, &
194 this%phmg_hrchy%lvl(i)%Xh)
195 call this%phmg_hrchy%lvl(i)%gs_h%init(this%phmg_hrchy%lvl(i)%dm_Xh)
196 call this%phmg_hrchy%lvl(i)%coef%init(this%phmg_hrchy%lvl(i)%gs_h)
199 do i = 0, this%nlvls - 1
200 call this%phmg_hrchy%lvl(i)%r%init(this%phmg_hrchy%lvl(i)%dm_Xh)
201 call this%phmg_hrchy%lvl(i)%w%init(this%phmg_hrchy%lvl(i)%dm_Xh)
202 call this%phmg_hrchy%lvl(i)%z%init(this%phmg_hrchy%lvl(i)%dm_Xh)
204 this%phmg_hrchy%lvl(i)%coef%ifh2 = coef%ifh2
205 call copy(this%phmg_hrchy%lvl(i)%coef%h1, coef%h1, &
206 this%phmg_hrchy%lvl(i)%dm_Xh%size())
208 call this%phmg_hrchy%lvl(i)%bc%init_base(this%phmg_hrchy%lvl(i)%coef)
209 if (bclst%size() .gt. 0 )
then
210 do j = 1, bclst%size()
212 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
215 call this%phmg_hrchy%lvl(i)%bc%finalize()
216 call this%phmg_hrchy%lvl(i)%bc%set_g(0.0_rp)
217 call this%phmg_hrchy%lvl(i)%bclst%init()
218 call this%phmg_hrchy%lvl(i)%bclst%append(this%phmg_hrchy%lvl(i)%bc)
221 if (trim(cheby_acc) .eq.
"schwarz")
then
222 call this%phmg_hrchy%lvl(i)%schwarz%init( &
223 this%phmg_hrchy%lvl(i)%Xh, &
224 this%phmg_hrchy%lvl(i)%dm_Xh, &
225 this%phmg_hrchy%lvl(i)%gs_h, &
226 this%phmg_hrchy%lvl(i)%bclst, &
231 call this%phmg_hrchy%lvl(i)%device_jacobi%init(&
232 this%phmg_hrchy%lvl(i)%coef, &
233 this%phmg_hrchy%lvl(i)%dm_Xh, &
234 this%phmg_hrchy%lvl(i)%gs_h)
236 call this%phmg_hrchy%lvl(i)%jacobi%init(&
237 this%phmg_hrchy%lvl(i)%coef, &
238 this%phmg_hrchy%lvl(i)%dm_Xh, &
239 this%phmg_hrchy%lvl(i)%gs_h)
243 if (trim(cheby_acc) .eq.
"jacobi")
then
244 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
245 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
246 this%phmg_hrchy%lvl(i)%device_jacobi)
249 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
250 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
252 if (trim(cheby_acc) .eq.
"schwarz")
then
253 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
254 this%phmg_hrchy%lvl(i)%schwarz
259 if (trim(cheby_acc) .eq.
"jacobi")
then
260 call this%phmg_hrchy%lvl(i)%cheby%init( &
261 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
262 this%phmg_hrchy%lvl(i)%jacobi)
265 call this%phmg_hrchy%lvl(i)%cheby%init( &
266 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
268 if (trim(cheby_acc) .eq.
"schwarz")
then
269 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
270 this%phmg_hrchy%lvl(i)%schwarz
281 call ax_helm_factory(this%ax, full_formulation = .false.)
284 allocate(this%intrp(this%nlvls - 1))
285 do i = 1, this%nlvls -1
286 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
287 this%phmg_hrchy%lvl(i)%Xh)
290 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
291 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
292 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
293 this%phmg_hrchy%lvl(this%nlvls -1)%bclst, &
294 crs_tamg_itrs, crs_tamg_cheby_degree)
299 class(
phmg_t),
intent(inout) :: this
303 class(
phmg_t),
intent(inout) :: this
304 integer,
intent(in) :: n
305 real(kind=
rp),
dimension(n),
intent(inout) :: z
306 real(kind=
rp),
dimension(n),
intent(inout) :: r
307 type(c_ptr) :: z_d, r_d
311 associate( mglvl => this%phmg_hrchy%lvl)
320 this%nlvls -1, mglvl, this%intrp, this%msh, this%Ax, &
323 call mglvl(0)%bclst%apply_scalar(mglvl(0)%z%x, n)
327 call copy(mglvl(0)%r%x, r, n)
329 mglvl(0)%z%x = 0.0_rp
330 mglvl(0)%w%x = 0.0_rp
333 this%nlvls -1, mglvl, this%intrp, this%msh, this%Ax, &
336 call mglvl(0)%bclst%apply_scalar(mglvl(0)%z%x, n)
337 call copy(z, mglvl(0)%z%x, n)
344 class(
phmg_t),
intent(inout) :: this
349 mg, intrp, msh, Ax, amg_solver)
350 type(ksp_monitor_t) :: ksp_results
353 type(interpolator_t) :: intrp(1:clvl)
354 type(tamg_solver_t),
intent(inout) :: amg_solver
355 class(ax_t),
intent(inout) :: ax
356 type(mesh_t),
intent(inout) :: msh
357 type(field_t) :: z, r, w
359 logical :: use_jacobi
363 call profiler_start_region(
'PHMG_cycle', 8)
367 call profiler_start_region(
'PHMG_PreSmooth', 9)
370 mg(lvl)%dm_Xh%size(), lvl)
372 if (neko_bcknd_device .eq. 1)
then
373 mg(lvl)%cheby_device%zero_initial_guess = .true.
374 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
375 r%x, mg(lvl)%dm_Xh%size(), &
376 mg(lvl)%coef, mg(lvl)%bclst, &
377 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
379 mg(lvl)%cheby%zero_initial_guess = .true.
380 ksp_results = mg(lvl)%cheby%solve(ax, z, &
381 r%x, mg(lvl)%dm_Xh%size(), &
382 mg(lvl)%coef, mg(lvl)%bclst, &
383 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
386 call profiler_end_region(
'PHMG_PreSmooth', 9)
391 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
392 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, glb_cmd_event)
393 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
394 call mg(lvl)%bclst%apply_scalar(w%x, mg(lvl)%dm_Xh%size())
396 if (neko_bcknd_device .eq. 1)
then
397 call device_sub3(w%x_d, r%x_d, w%x_d, mg(lvl)%dm_Xh%size())
405 if (neko_bcknd_device .eq. 1)
then
406 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
408 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
411 call profiler_start_region(
'PHMG_map_to_coarse', 9)
412 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
413 call profiler_end_region(
'PHMG_map_to_coarse', 9)
415 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
416 gs_op_add, glb_cmd_event)
417 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
419 call mg(lvl+1)%bclst%apply_scalar( &
421 mg(lvl+1)%dm_Xh%size())
425 if (neko_bcknd_device .eq. 1)
then
426 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
428 mg(lvl+1)%z%x = 0.0_rp
430 if (lvl+1 .eq. clvl)
then
431 call profiler_start_region(
'PHMG_tAMG_coarse_grid', 9)
432 call amg_solver%solve(mg(lvl+1)%z%x, &
434 mg(lvl+1)%dm_Xh%size())
435 call profiler_end_region(
'PHMG_tAMG_coarse_grid', 9)
438 call phmg_mg_cycle(mg(lvl+1)%z, mg(lvl+1)%r, mg(lvl+1)%w, lvl+1, &
439 clvl, mg, intrp, msh, ax, amg_solver)
445 call profiler_start_region(
'PHMG_map_to_fine', 9)
446 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
447 call profiler_end_region(
'PHMG_map_to_fine', 9)
449 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, glb_cmd_event)
450 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
452 if (neko_bcknd_device .eq. 1)
then
453 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
455 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
461 if (neko_bcknd_device .eq. 1)
then
462 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
470 call profiler_start_region(
'PHMG_PostSmooth', 9)
473 mg(lvl)%dm_Xh%size(), lvl)
475 if (neko_bcknd_device .eq. 1)
then
476 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
477 r%x, mg(lvl)%dm_Xh%size(), &
478 mg(lvl)%coef, mg(lvl)%bclst, &
479 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
481 ksp_results = mg(lvl)%cheby%solve(ax, z, &
482 r%x, mg(lvl)%dm_Xh%size(), &
483 mg(lvl)%coef, mg(lvl)%bclst, &
484 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
487 call profiler_end_region(
'PHMG_PostSmooth', 9)
489 call profiler_end_region(
'PHMG_cycle', 8)
503 class(ax_t),
intent(inout) :: Ax
504 type(mesh_t),
intent(inout) :: msh
505 type(field_t),
intent(inout) :: z, r, w
506 integer,
intent(in) :: n, lvl
507 integer :: i, iblk, ni, niblk
510 if (neko_bcknd_device .eq. 1)
then
512 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
513 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
514 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
515 call mg%bclst%apply_scalar(w%x, n)
516 call device_sub3(w%x_d, r%x_d, w%x_d, n)
518 call mg%device_jacobi%solve(w%x, w%x, n)
520 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
524 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
525 call mg%gs_h%op(w%x, n, gs_op_add)
526 call mg%bclst%apply_scalar(w%x, n)
527 call sub3(w%x, r%x, w%x, n)
529 call mg%jacobi%solve(w%x, w%x, n)
531 call add2s2(z%x, w%x, 0.6_rp, n)
540 class(ax_t),
intent(inout) :: Ax
541 type(mesh_t),
intent(inout) :: msh
542 type(field_t) :: z, r, w
544 character(len=LOG_SIZE) :: log_buf
545 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
546 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
547 call mg%bclst%apply_scalar(w%x, mg%dm_Xh%size())
548 call device_sub3(w%x_d, r%x_d, w%x_d, mg%dm_Xh%size())
549 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
551 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO - PRE', lvl, val
552 else if (typ .eq. 2)
then
553 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO -POST', lvl, val
554 else if (typ .eq. 3)
then
555 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO- PRE', lvl, val
556 else if (typ .eq. 4)
then
557 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO-POST', lvl, val
558 else if (typ .eq. 5)
then
559 write(log_buf,
'(A15,I4,F12.6)')
'TAMG - PRE', lvl, val
560 else if (typ .eq. 6)
then
561 write(log_buf,
'(A15,I4,F12.6)')
'TAMG -POST', lvl, val
563 write(log_buf,
'(A15,I4,F12.6)')
'RESID', lvl, val
565 call neko_log%message(log_buf)
569 integer,
intent(in) :: nlvls
570 integer,
intent(in) :: smoo_type
573 character(len=LOG_SIZE) :: log_buf, smoo_name
575 call neko_log%section(
'PHMG')
577 if (smoo_type .eq. 1)
then
578 write(smoo_name,
'(A16)')
'CHEBY-acc JACOBI'
579 else if (smoo_type .eq. 2)
then
580 write(smoo_name,
'(A17)')
'CHEBY-acc SCHWARZ'
582 write(smoo_name,
'(A5)')
'CHEBY'
585 write(log_buf,
'(A28,I2,A8)') &
586 'Creating PHMG hierarchy with', &
588 call neko_log%message(log_buf)
592 write(log_buf,
'(A8,I2,A8,I2)') &
593 '-- level', i,
'-- lx:',
phmg%lvl(i)%Xh%lx
594 call neko_log%message(log_buf)
596 if (i .eq. clvl)
then
597 write(log_buf,
'(A19,A20)') &
599 call neko_log%message(log_buf)
601 write(log_buf,
'(A22,A20)') &
604 call neko_log%message(log_buf)
606 write(log_buf,
'(A28,I2)') &
608 phmg%lvl(i)%smoother_itrs
609 call neko_log%message(log_buf)
613 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...
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_solve(this, z, r, n)
subroutine phmg_init_from_components(this, coef, bclst, smoother_itrs, cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
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.