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 => &
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(:)
136 crs_tamg_cheby_degree, 4)
138 if (phmg_params%valid_path(
'pcoarsening_schedule'))
then
139 call json_get(phmg_params,
'pcoarsening_schedule', pcrs_sched)
141 allocate(pcrs_sched(2))
147 call this%init_from_components(coef, bclst, smoother_itrs, &
148 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree,&
154 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, &
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
174 this%nlvls =
size(pcrs_sched) + 1
175 allocate(lx_lvls(0:this%nlvls - 1))
176 lx_lvls(1:) = pcrs_sched + 1
178 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
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
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)
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), &
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)
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)
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())
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()
216 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
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)
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, &
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)
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)
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)
253 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
254 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
256 if (trim(cheby_acc) .eq.
"schwarz")
then
257 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
258 this%phmg_hrchy%lvl(i)%schwarz
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)
269 call this%phmg_hrchy%lvl(i)%cheby%init( &
270 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
272 if (trim(cheby_acc) .eq.
"schwarz")
then
273 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
274 this%phmg_hrchy%lvl(i)%schwarz
285 call ax_helm_factory(this%ax, full_formulation = .false.)
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)
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)
303 class(
phmg_t),
intent(inout) :: this
307 class(
phmg_t),
intent(inout) :: this
308 integer,
intent(in) :: n
309 real(kind=
rp),
dimension(n),
intent(inout) :: z
310 real(kind=
rp),
dimension(n),
intent(inout) :: r
311 type(c_ptr) :: z_d, r_d
316 associate( mglvl => this%phmg_hrchy%lvl)
336 mglvl(0)%r%x(i,1,1,1) = r(i)
338 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
339 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
350 z(i) = mglvl(0)%z%x(i,1,1,1)
360 class(
phmg_t),
intent(inout) :: this
365 class(
phmg_t),
intent(inout) :: this
366 type(ksp_monitor_t) :: ksp_results
367 character(len=2) :: lvl_name
370 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
371 msh => this%msh, ax => this%Ax)
372 do lvl = 0, this%nlvls-2
373 write(lvl_name,
'(I0)') lvl
374 call profiler_start_region(
"PHMG_level_" // trim(lvl_name))
375 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
379 if (neko_bcknd_device .eq. 1)
then
380 mg(lvl)%cheby_device%zero_initial_guess = .true.
381 ksp_results = mg(lvl)%cheby_device%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)
386 mg(lvl)%cheby%zero_initial_guess = .true.
387 ksp_results = mg(lvl)%cheby%solve(ax, z, &
388 r%x, mg(lvl)%dm_Xh%size(), &
389 mg(lvl)%coef, mg(lvl)%bclst, &
390 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
396 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
397 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
399 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
400 call mg(lvl)%bclst%apply_scalar(w%x, mg(lvl)%dm_Xh%size())
402 if (neko_bcknd_device .eq. 1)
then
403 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
409 do i = 1, mg(lvl)%dm_Xh%size()
410 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
418 if (neko_bcknd_device .eq. 1)
then
419 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
421 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
424 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
426 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
427 gs_op_add, glb_cmd_event)
428 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
430 call mg(lvl+1)%bclst%apply_scalar( &
432 mg(lvl+1)%dm_Xh%size())
434 if (neko_bcknd_device .eq. 1)
then
435 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
441 do i = 1, mg(lvl+1)%dm_Xh%size()
442 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
447 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
450 call profiler_start_region(
'PHMG_coarse-solve' )
454 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
455 mg(this%nlvls-1)%r%x, &
456 mg(this%nlvls-1)%dm_Xh%size())
457 call profiler_end_region(
'PHMG_coarse-solve' )
459 do lvl = (this%nlvls-2), 0, -1
460 write(lvl_name,
'(I0)') lvl
461 call profiler_start_region(
"PHMG_level_" // trim(lvl_name))
462 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
466 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
468 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
470 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
472 if (neko_bcknd_device .eq. 1)
then
473 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
475 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
481 if (neko_bcknd_device .eq. 1)
then
482 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
488 do i = 1, mg(lvl)%dm_Xh%size()
489 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
497 if (neko_bcknd_device .eq. 1)
then
498 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
499 r%x, mg(lvl)%dm_Xh%size(), &
500 mg(lvl)%coef, mg(lvl)%bclst, &
501 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
503 ksp_results = mg(lvl)%cheby%solve(ax, z, &
504 r%x, mg(lvl)%dm_Xh%size(), &
505 mg(lvl)%coef, mg(lvl)%bclst, &
506 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
509 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
526 class(ax_t),
intent(inout) :: Ax
527 type(mesh_t),
intent(inout) :: msh
528 type(field_t),
intent(inout) :: z, r, w
529 integer,
intent(in) :: n, lvl
530 integer :: i, j, iblk, ni, niblk
532 ni = mg%smoother_itrs
533 if (neko_bcknd_device .eq. 1)
then
535 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
536 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
537 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
538 call mg%bclst%apply_scalar(w%x, n)
539 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
541 call mg%device_jacobi%solve(w%x, w%x, n)
543 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
547 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
548 call mg%gs_h%op(w%x, n, gs_op_add)
549 call mg%bclst%apply_scalar(w%x, n)
550 call add2s1(w%x, r%x, -1.0_rp, n)
552 call mg%jacobi%solve(w%x, w%x, n)
554 call add2s2(z%x, w%x, 0.6_rp, n)
563 class(ax_t),
intent(inout) :: Ax
564 type(mesh_t),
intent(inout) :: msh
565 type(field_t) :: z, r, w
567 character(len=LOG_SIZE) :: log_buf
568 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
569 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
570 call mg%bclst%apply_scalar(w%x, mg%dm_Xh%size())
571 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
572 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
574 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO - PRE', lvl, val
575 else if (typ .eq. 2)
then
576 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO -POST', lvl, val
577 else if (typ .eq. 3)
then
578 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO- PRE', lvl, val
579 else if (typ .eq. 4)
then
580 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO-POST', lvl, val
581 else if (typ .eq. 5)
then
582 write(log_buf,
'(A15,I4,F12.6)')
'TAMG - PRE', lvl, val
583 else if (typ .eq. 6)
then
584 write(log_buf,
'(A15,I4,F12.6)')
'TAMG -POST', lvl, val
586 write(log_buf,
'(A15,I4,F12.6)')
'RESID', lvl, val
588 call neko_log%message(log_buf)
592 integer,
intent(in) :: nlvls
593 integer,
intent(in) :: smoo_type
596 character(len=LOG_SIZE) :: log_buf, smoo_name
598 call neko_log%section(
'PHMG')
600 if (smoo_type .eq. 1)
then
601 write(smoo_name,
'(A16)')
'CHEBY-acc JACOBI'
602 else if (smoo_type .eq. 2)
then
603 write(smoo_name,
'(A17)')
'CHEBY-acc SCHWARZ'
605 write(smoo_name,
'(A5)')
'CHEBY'
608 write(log_buf,
'(A28,I2,A8)') &
609 'Creating PHMG hierarchy with', &
611 call neko_log%message(log_buf)
615 write(log_buf,
'(A8,I2,A8,I2)') &
616 '-- level', i,
'-- lx:',
phmg%lvl(i)%Xh%lx
617 call neko_log%message(log_buf)
619 if (i .eq. clvl)
then
620 write(log_buf,
'(A19,A20)') &
622 call neko_log%message(log_buf)
624 write(log_buf,
'(A22,A20)') &
627 call neko_log%message(log_buf)
629 write(log_buf,
'(A28,I2)') &
631 phmg%lvl(i)%smoother_itrs
632 call neko_log%message(log_buf)
636 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_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.
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 add2s1(a, b, c1, n)
Vector addition with scalar multiplication (multiplication on first argument)
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_mg_cycle(this)
subroutine phmg_resid_monitor(z, r, w, mg, msh, ax, lvl, typ)
subroutine print_phmg_info(nlvls, smoo_type, phmg)
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.