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
306 call this%amg_solver%free()
308 if (
allocated(this%intrp))
then
309 do i = 1,
size(this%intrp)
310 call this%intrp(i)%free()
312 deallocate(this%intrp)
315 if (
allocated(this%ax))
then
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()
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()
330 if (
allocated(this%phmg_hrchy%lvl(i)%schwarz%work1))
then
331 call this%phmg_hrchy%lvl(i)%schwarz%free()
334 call this%phmg_hrchy%lvl(i)%bclst%free()
335 call this%phmg_hrchy%lvl(i)%bc%free()
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)
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)
355 deallocate(this%phmg_hrchy%lvl)
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
372 associate( mglvl => this%phmg_hrchy%lvl)
393 mglvl(0)%r%x(i,1,1,1) = r(i)
395 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
396 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
408 z(i) = mglvl(0)%z%x(i,1,1,1)
418 class(
phmg_t),
intent(inout) :: this
423 class(
phmg_t),
intent(inout) :: this
424 type(ksp_monitor_t) :: ksp_results
425 character(len=2) :: lvl_name
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)
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)
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)
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, &
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())
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())
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)
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())
480 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
483 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
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)
489 call mg(lvl+1)%bclst%apply_scalar( &
491 mg(lvl+1)%dm_Xh%size())
493 if (neko_bcknd_device .eq. 1)
then
494 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
501 do i = 1, mg(lvl+1)%dm_Xh%size()
502 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
507 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
510 call profiler_start_region(
'PHMG_coarse-solve' )
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' )
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)
526 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
528 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
530 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
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())
535 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
541 if (neko_bcknd_device .eq. 1)
then
542 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
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)
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)
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)
570 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
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
593 ni = mg%smoother_itrs
594 if (neko_bcknd_device .eq. 1)
then
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)
602 call mg%device_jacobi%solve(w%x, w%x, n)
604 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
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)
613 call mg%jacobi%solve(w%x, w%x, n)
615 call add2s2(z%x, w%x, 0.6_rp, n)
624 class(ax_t),
intent(inout) :: Ax
625 type(mesh_t),
intent(inout) :: msh
626 type(field_t) :: z, r, w
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())
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
647 write(log_buf,
'(A15,I4,F12.6)')
'RESID', lvl, val
649 call neko_log%message(log_buf)
653 integer,
intent(in) :: nlvls
654 integer,
intent(in) :: smoo_type
657 character(len=LOG_SIZE) :: log_buf, smoo_name
659 call neko_log%section(
'PHMG')
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'
666 write(smoo_name,
'(A5)')
'CHEBY'
669 write(log_buf,
'(A28,I2,A8)') &
670 'Creating PHMG hierarchy with', &
672 call neko_log%message(log_buf)
676 write(log_buf,
'(A8,I2,A8,I2)') &
677 '-- level', i,
'-- lx:',
phmg%lvl(i)%Xh%lx
678 call neko_log%message(log_buf)
680 if (i .eq. clvl)
then
681 write(log_buf,
'(A19,A20)') &
683 call neko_log%message(log_buf)
685 write(log_buf,
'(A22,A20)') &
688 call neko_log%message(log_buf)
690 write(log_buf,
'(A28,I2)') &
692 phmg%lvl(i)%smoother_itrs
693 call neko_log%message(log_buf)
697 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.