56 use json_module,
only : json_file
68 use,
intrinsic :: iso_c_binding
75 integer :: smoother_itrs = 10
78 type(
gs_t),
pointer :: gs_h
99 class(
ax_t),
allocatable :: ax
105 integer :: last_metrics_version = -1
109 logical :: update_enabled = .false.
111 logical :: refresh_eigs = .true.
113 integer :: refresh_eigs_frequency = 20
115 logical :: eigs_warm_start = .true.
117 integer :: power_its_refresh = 20
119 integer :: n_refresh = 0
121 character(len=:),
allocatable :: cheby_acc
124 procedure, pass(this) :: init_from_components => &
135 class(
phmg_t),
intent(inout),
target :: this
136 type(
coef_t),
intent(in),
target :: coef
137 type(
bc_list_t),
intent(inout),
target :: bclst
138 type(json_file),
intent(inout) :: phmg_params
139 integer :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
140 integer :: smoother_itrs
141 character(len=:),
allocatable :: cheby_acc
142 integer,
allocatable :: pcrs_sched(:)
143 logical :: update_enabled
158 crs_tamg_cheby_degree, 4)
160 if (phmg_params%valid_path(
'pcoarsening_schedule'))
then
161 call json_get(phmg_params,
'pcoarsening_schedule', pcrs_sched)
163 allocate(pcrs_sched(2))
169 update_enabled, .false.)
173 this%refresh_eigs, .true.)
175 this%refresh_eigs_frequency, 20)
177 this%eigs_warm_start, .true.)
179 this%power_its_refresh, 20)
181 call this%init_from_components(coef, bclst, smoother_itrs, &
182 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree,&
183 pcrs_sched, update_enabled)
188 cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, &
189 pcrs_sched, update_enabled)
190 class(
phmg_t),
intent(inout),
target :: this
191 type(
coef_t),
intent(in),
target :: coef
192 type(
bc_list_t),
intent(inout),
target :: bclst
193 integer,
intent(in) :: smoother_itrs
194 character(len=:),
allocatable :: cheby_acc
195 integer,
intent(in) :: crs_tamg_lvls, crs_tamg_itrs
196 integer,
intent(in) :: crs_tamg_cheby_degree
197 integer,
intent(in),
allocatable :: pcrs_sched(:)
198 logical,
intent(in),
optional :: update_enabled
199 integer :: lx_crs, lx_mid
200 integer,
allocatable :: lx_lvls(:)
201 integer :: n, i, j, st
202 class(
bc_t),
pointer :: bc_j
203 logical :: use_jacobi, use_cheby
208 this%cheby_acc = trim(cheby_acc)
210 if (
present(update_enabled))
then
211 this%update_enabled = update_enabled
213 this%update_enabled = .false.
216 this%nlvls =
size(pcrs_sched) + 1
217 allocate(lx_lvls(0:this%nlvls - 1))
218 lx_lvls(1:) = pcrs_sched + 1
220 allocate(this%phmg_hrchy%lvl(0:this%nlvls - 1))
222 this%phmg_hrchy%lvl(0)%lvl = 0
223 this%phmg_hrchy%lvl(0)%smoother_itrs = smoother_itrs
224 this%phmg_hrchy%lvl(0)%Xh => coef%Xh
225 this%phmg_hrchy%lvl(0)%coef => coef
226 this%phmg_hrchy%lvl(0)%dm_Xh => coef%dof
227 this%phmg_hrchy%lvl(0)%gs_h => coef%gs_h
229 do i = 1, this%nlvls - 1
230 allocate(this%phmg_hrchy%lvl(i)%Xh)
231 allocate(this%phmg_hrchy%lvl(i)%dm_Xh)
232 allocate(this%phmg_hrchy%lvl(i)%gs_h)
233 allocate(this%phmg_hrchy%lvl(i)%coef)
235 this%phmg_hrchy%lvl(i)%lvl = i
236 this%phmg_hrchy%lvl(i)%smoother_itrs = smoother_itrs
237 call this%phmg_hrchy%lvl(i)%Xh%init(
gll, lx_lvls(i), lx_lvls(i), &
239 call this%phmg_hrchy%lvl(i)%dm_Xh%init(coef%msh, &
240 this%phmg_hrchy%lvl(i)%Xh)
241 call this%phmg_hrchy%lvl(i)%gs_h%init(this%phmg_hrchy%lvl(i)%dm_Xh)
246 if (this%update_enabled)
then
247 call this%phmg_hrchy%lvl(i)%coef%init( &
250 call this%phmg_hrchy%lvl(i)%coef%init( &
255 do i = 0, this%nlvls - 1
256 call this%phmg_hrchy%lvl(i)%r%init(this%phmg_hrchy%lvl(i)%dm_Xh)
257 call this%phmg_hrchy%lvl(i)%w%init(this%phmg_hrchy%lvl(i)%dm_Xh)
258 call this%phmg_hrchy%lvl(i)%z%init(this%phmg_hrchy%lvl(i)%dm_Xh)
260 this%phmg_hrchy%lvl(i)%coef%ifh2 = coef%ifh2
261 call copy(this%phmg_hrchy%lvl(i)%coef%h1, coef%h1, &
262 this%phmg_hrchy%lvl(i)%dm_Xh%size())
264 call this%phmg_hrchy%lvl(i)%bc%init_base(this%phmg_hrchy%lvl(i)%coef)
265 if (bclst%size() .gt. 0 )
then
266 do j = 1, bclst%size()
268 call this%phmg_hrchy%lvl(i)%bc%mark_facets(bc_j%marked_facet)
271 call this%phmg_hrchy%lvl(i)%bc%finalize()
272 call this%phmg_hrchy%lvl(i)%bc%set_g(0.0_rp)
273 call this%phmg_hrchy%lvl(i)%bc_projector%mark(this%phmg_hrchy%lvl(i)%bc)
276 if (trim(cheby_acc) .eq.
"schwarz")
then
277 call this%phmg_hrchy%lvl(i)%schwarz%init( &
278 this%phmg_hrchy%lvl(i)%Xh, &
279 this%phmg_hrchy%lvl(i)%dm_Xh, &
280 this%phmg_hrchy%lvl(i)%gs_h, &
281 this%phmg_hrchy%lvl(i)%bc_projector, &
286 call this%phmg_hrchy%lvl(i)%device_jacobi%init(&
287 this%phmg_hrchy%lvl(i)%coef, &
288 this%phmg_hrchy%lvl(i)%dm_Xh, &
289 this%phmg_hrchy%lvl(i)%gs_h)
291 call this%phmg_hrchy%lvl(i)%jacobi%init(&
292 this%phmg_hrchy%lvl(i)%coef, &
293 this%phmg_hrchy%lvl(i)%dm_Xh, &
294 this%phmg_hrchy%lvl(i)%gs_h)
298 if (trim(cheby_acc) .eq.
"jacobi")
then
299 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
300 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
301 this%phmg_hrchy%lvl(i)%device_jacobi)
304 call this%phmg_hrchy%lvl(i)%cheby_device%init( &
305 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
307 if (trim(cheby_acc) .eq.
"schwarz")
then
308 this%phmg_hrchy%lvl(i)%cheby_device%schwarz => &
309 this%phmg_hrchy%lvl(i)%schwarz
314 if (trim(cheby_acc) .eq.
"jacobi")
then
315 call this%phmg_hrchy%lvl(i)%cheby%init( &
316 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs, &
317 this%phmg_hrchy%lvl(i)%jacobi)
320 call this%phmg_hrchy%lvl(i)%cheby%init( &
321 this%phmg_hrchy%lvl(i)%dm_Xh%size(), smoother_itrs)
323 if (trim(cheby_acc) .eq.
"schwarz")
then
324 this%phmg_hrchy%lvl(i)%cheby%schwarz => &
325 this%phmg_hrchy%lvl(i)%schwarz
336 call ax_helm_allocator(this%ax, type_name =
"standard")
339 allocate(this%intrp(this%nlvls - 1))
340 do i = 1, this%nlvls -1
341 call this%intrp(i)%init(this%phmg_hrchy%lvl(i-1)%Xh, &
342 this%phmg_hrchy%lvl(i)%Xh)
346 if (this%update_enabled)
then
347 allocate(this%crd_intrp(this%nlvls - 1))
348 do i = 1, this%nlvls - 1
349 call this%crd_intrp(i)%init(this%phmg_hrchy%lvl(i)%Xh, &
350 this%phmg_hrchy%lvl(0)%Xh)
354 call this%amg_solver%init(this%ax, this%phmg_hrchy%lvl(this%nlvls -1)%Xh, &
355 this%phmg_hrchy%lvl(this%nlvls -1)%coef, this%msh, &
356 this%phmg_hrchy%lvl(this%nlvls-1)%gs_h, crs_tamg_lvls, &
357 this%phmg_hrchy%lvl(this%nlvls -1)%bc_projector, &
358 crs_tamg_itrs, crs_tamg_cheby_degree)
361 this%last_metrics_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
364 if (this%update_enabled)
then
365 do i = 0, this%nlvls - 1
366 this%phmg_hrchy%lvl(i)%cheby%warm_start_eigs = this%eigs_warm_start
367 this%phmg_hrchy%lvl(i)%cheby%power_its_refresh = &
368 this%power_its_refresh
369 this%phmg_hrchy%lvl(i)%cheby_device%warm_start_eigs = &
371 this%phmg_hrchy%lvl(i)%cheby_device%power_its_refresh = &
372 this%power_its_refresh
374 call this%amg_solver%set_eig_refresh(this%eigs_warm_start, &
375 this%power_its_refresh)
377 if (trim(cheby_acc) .eq.
"schwarz")
then
378 call neko_warning(
"PHMG: the Schwarz smoother is not refreshed " // &
379 "when the mesh changes. Its local solves stay at the " // &
387 class(
phmg_t),
intent(inout) :: this
390 call this%amg_solver%free()
392 if (
allocated(this%intrp))
then
393 do i = 1,
size(this%intrp)
394 call this%intrp(i)%free()
396 deallocate(this%intrp)
399 if (
allocated(this%crd_intrp))
then
400 do i = 1,
size(this%crd_intrp)
401 call this%crd_intrp(i)%free()
403 deallocate(this%crd_intrp)
406 if (
allocated(this%ax))
then
411 if (
allocated(this%phmg_hrchy%lvl))
then
412 do i = lbound(this%phmg_hrchy%lvl, 1), ubound(this%phmg_hrchy%lvl, 1)
413 call this%phmg_hrchy%lvl(i)%r%free()
414 call this%phmg_hrchy%lvl(i)%w%free()
415 call this%phmg_hrchy%lvl(i)%z%free()
417 call this%phmg_hrchy%lvl(i)%cheby%free()
418 call this%phmg_hrchy%lvl(i)%cheby_device%free()
419 call this%phmg_hrchy%lvl(i)%jacobi%free()
420 call this%phmg_hrchy%lvl(i)%device_jacobi%free()
422 if (
allocated(this%phmg_hrchy%lvl(i)%schwarz%work1))
then
423 call this%phmg_hrchy%lvl(i)%schwarz%free()
426 call this%phmg_hrchy%lvl(i)%bc_projector%free()
427 call this%phmg_hrchy%lvl(i)%bc%free()
431 if (this%phmg_hrchy%lvl(i)%lvl .gt. 0)
then
432 call this%phmg_hrchy%lvl(i)%coef%free()
433 call this%phmg_hrchy%lvl(i)%gs_h%free()
434 call this%phmg_hrchy%lvl(i)%dm_Xh%free()
435 call this%phmg_hrchy%lvl(i)%Xh%free()
436 deallocate(this%phmg_hrchy%lvl(i)%coef)
437 deallocate(this%phmg_hrchy%lvl(i)%gs_h)
438 deallocate(this%phmg_hrchy%lvl(i)%dm_Xh)
439 deallocate(this%phmg_hrchy%lvl(i)%Xh)
442 nullify(this%phmg_hrchy%lvl(i)%coef)
443 nullify(this%phmg_hrchy%lvl(i)%gs_h)
444 nullify(this%phmg_hrchy%lvl(i)%dm_Xh)
445 nullify(this%phmg_hrchy%lvl(i)%Xh)
447 deallocate(this%phmg_hrchy%lvl)
455 class(
phmg_t),
intent(inout) :: this
456 integer,
intent(in) :: n
457 real(kind=
rp),
dimension(n),
intent(inout) :: z
458 real(kind=
rp),
dimension(n),
intent(inout) :: r
459 type(c_ptr) :: z_d, r_d
464 associate( mglvl => this%phmg_hrchy%lvl)
485 mglvl(0)%r%x(i,1,1,1) = r(i)
487 mglvl(0)%z%x(i,1,1,1) = 0.0_rp
488 mglvl(0)%w%x(i,1,1,1) = 0.0_rp
500 z(i) = mglvl(0)%z%x(i,1,1,1)
511 class(
phmg_t),
intent(inout) :: this
512 integer :: fine_version
515 if (.not. this%update_enabled)
return
517 fine_version = this%phmg_hrchy%lvl(0)%coef%metrics_version
520 if (fine_version .eq. this%last_metrics_version)
return
526 do_eigs = (this%refresh_eigs) .and. &
527 (this%refresh_eigs_frequency .gt. 0) .and. &
528 (mod(this%n_refresh,
max(this%refresh_eigs_frequency, 1)) &
532 this%n_refresh = this%n_refresh + 1
533 this%last_metrics_version = fine_version
541 class(
phmg_t),
intent(inout) :: this
544 if (.not.
allocated(this%crd_intrp))
then
545 call neko_error(
"PHMG: update requested but coordinate " // &
546 "interpolators were not initialized.")
549 call profiler_start_region(
'PHMG_update_geometry')
550 associate(mg => this%phmg_hrchy%lvl, nelv => this%msh%nelv)
551 do i = 1, this%nlvls - 1
552 call this%crd_intrp(i)%map(mg(i)%dm_Xh%x%x, mg(0)%dm_Xh%x%x, &
554 call this%crd_intrp(i)%map(mg(i)%dm_Xh%y%x, mg(0)%dm_Xh%y%x, &
556 call this%crd_intrp(i)%map(mg(i)%dm_Xh%z%x, mg(0)%dm_Xh%z%x, &
559 call mg(i)%coef%recompute_metrics()
562 call profiler_end_region(
'PHMG_update_geometry')
570 class(
phmg_t),
intent(inout) :: this
573 call profiler_start_region(
'PHMG_update_smoother_acc')
574 associate(mg => this%phmg_hrchy%lvl)
575 select case (trim(this%cheby_acc))
577 do i = 0, this%nlvls - 1
578 if (neko_bcknd_device .eq. 1)
then
579 call mg(i)%device_jacobi%update()
581 call mg(i)%jacobi%update()
591 call profiler_end_region(
'PHMG_update_smoother_acc')
598 class(
phmg_t),
intent(inout) :: this
601 associate(mg => this%phmg_hrchy%lvl)
602 do i = 0, this%nlvls - 1
603 if (neko_bcknd_device .eq. 1)
then
604 mg(i)%cheby_device%recompute_eigs = .true.
606 mg(i)%cheby%recompute_eigs = .true.
611 call this%amg_solver%invalidate_eigs()
617 class(
phmg_t),
intent(inout) :: this
618 type(ksp_monitor_t) :: ksp_results
619 character(len=2) :: lvl_name
622 associate(mg => this%phmg_hrchy%lvl, intrp => this%intrp, &
623 msh => this%msh, ax => this%Ax)
624 do lvl = 0, this%nlvls-2
625 write(lvl_name,
'(I0)') lvl
626 call profiler_start_region(
"PHMG_level_" // trim(lvl_name))
627 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
631 if (neko_bcknd_device .eq. 1)
then
632 mg(lvl)%cheby_device%zero_initial_guess = .true.
633 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
634 r%x, mg(lvl)%dm_Xh%size(), &
635 mg(lvl)%coef, mg(lvl)%bc_projector, &
636 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
638 mg(lvl)%cheby%zero_initial_guess = .true.
639 ksp_results = mg(lvl)%cheby%solve(ax, z, &
640 r%x, mg(lvl)%dm_Xh%size(), &
641 mg(lvl)%coef, mg(lvl)%bc_projector, &
642 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
648 call ax%compute(w%x, z%x, mg(lvl)%coef, msh, mg(lvl)%Xh)
649 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
651 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
652 call mg(lvl)%bc_projector%apply(w%x, mg(lvl)%dm_Xh%size())
654 if (neko_bcknd_device .eq. 1)
then
655 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg(lvl)%dm_Xh%size())
662 do i = 1, mg(lvl)%dm_Xh%size()
663 w%x(i,1,1,1) = r%x(i,1,1,1) - w%x(i,1,1,1)
671 if (neko_bcknd_device .eq. 1)
then
672 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
674 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
677 call intrp(lvl+1)%map(mg(lvl+1)%r%x, w%x, msh%nelv, mg(lvl+1)%Xh)
679 call mg(lvl+1)%gs_h%op(mg(lvl+1)%r%x, mg(lvl+1)%dm_Xh%size(), &
680 gs_op_add, glb_cmd_event)
681 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
683 call mg(lvl+1)%bc_projector%apply( &
685 mg(lvl+1)%dm_Xh%size())
687 if (neko_bcknd_device .eq. 1)
then
688 call device_rzero(mg(lvl+1)%z%x_d, mg(lvl+1)%dm_Xh%size())
695 do i = 1, mg(lvl+1)%dm_Xh%size()
696 mg(lvl+1)%z%x(i,1,1,1) = 0.0_rp
701 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
704 call profiler_start_region(
'PHMG_coarse-solve' )
708 call this%amg_solver%solve(mg(this%nlvls-1)%z%x, &
709 mg(this%nlvls-1)%r%x, &
710 mg(this%nlvls-1)%dm_Xh%size())
711 call profiler_end_region(
'PHMG_coarse-solve' )
713 do lvl = (this%nlvls-2), 0, -1
714 write(lvl_name,
'(I0)') lvl
715 call profiler_start_region(
"PHMG_level_" // trim(lvl_name))
716 associate(z => mg(lvl)%z, r => mg(lvl)%r, w => mg(lvl)%w)
720 call intrp(lvl+1)%map(w%x, mg(lvl+1)%z%x, msh%nelv, mg(lvl)%Xh)
722 call mg(lvl)%gs_h%op(w%x, mg(lvl)%dm_Xh%size(), gs_op_add, &
724 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
726 if (neko_bcknd_device .eq. 1)
then
727 call device_col2(w%x_d, mg(lvl)%coef%mult_d, mg(lvl)%dm_Xh%size())
729 call col2(w%x, mg(lvl)%coef%mult, mg(lvl)%dm_Xh%size())
735 if (neko_bcknd_device .eq. 1)
then
736 call device_add2(z%x_d, w%x_d, mg(lvl)%dm_Xh%size())
743 do i = 1, mg(lvl)%dm_Xh%size()
744 z%x(i,1,1,1) = z%x(i,1,1,1) + w%x(i,1,1,1)
752 if (neko_bcknd_device .eq. 1)
then
753 ksp_results = mg(lvl)%cheby_device%solve(ax, z, &
754 r%x, mg(lvl)%dm_Xh%size(), &
755 mg(lvl)%coef, mg(lvl)%bc_projector, &
756 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
758 ksp_results = mg(lvl)%cheby%solve(ax, z, &
759 r%x, mg(lvl)%dm_Xh%size(), &
760 mg(lvl)%coef, mg(lvl)%bc_projector, &
761 mg(lvl)%gs_h, niter = mg(lvl)%smoother_itrs)
764 call profiler_end_region(
"PHMG_level_" // trim(lvl_name))
781 class(ax_t),
intent(inout) :: Ax
782 type(mesh_t),
intent(inout) :: msh
783 type(field_t),
intent(inout) :: z, r, w
784 integer,
intent(in) :: n, lvl
785 integer :: i, j, iblk, ni, niblk
787 ni = mg%smoother_itrs
788 if (neko_bcknd_device .eq. 1)
then
790 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
791 call mg%gs_h%op(w%x, n, gs_op_add, glb_cmd_event)
792 call device_stream_wait_event(glb_cmd_queue, glb_cmd_event, 0)
793 call mg%bc_projector%apply(w%x, n)
794 call device_add2s1(w%x_d, r%x_d, -1.0_rp, n)
796 call mg%device_jacobi%solve(w%x, w%x, n)
798 call device_add2s2(z%x_d, w%x_d, 0.6_rp, n)
802 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
803 call mg%gs_h%op(w%x, n, gs_op_add)
804 call mg%bc_projector%apply(w%x, n)
805 call add2s1(w%x, r%x, -1.0_rp, n)
807 call mg%jacobi%solve(w%x, w%x, n)
809 call add2s2(z%x, w%x, 0.6_rp, n)
818 class(ax_t),
intent(inout) :: Ax
819 type(mesh_t),
intent(inout) :: msh
820 type(field_t) :: z, r, w
822 character(len=LOG_SIZE) :: log_buf
823 call ax%compute(w%x, z%x, mg%coef, msh, mg%Xh)
824 call mg%gs_h%op(w%x, mg%dm_Xh%size(), gs_op_add)
825 call mg%bc_projector%apply(w%x, mg%dm_Xh%size())
826 call device_add2s1(w%x_d, r%x_d, -1.0_rp, mg%dm_Xh%size())
827 val = device_glsc2(w%x_d, w%x_d, mg%dm_Xh%size())
829 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO - PRE', lvl, val
830 else if (typ .eq. 2)
then
831 write(log_buf,
'(A15,I4,F12.6)')
'PRESMOO -POST', lvl, val
832 else if (typ .eq. 3)
then
833 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO- PRE', lvl, val
834 else if (typ .eq. 4)
then
835 write(log_buf,
'(A15,I4,F12.6)')
'POSTSMOO-POST', lvl, val
836 else if (typ .eq. 5)
then
837 write(log_buf,
'(A15,I4,F12.6)')
'TAMG - PRE', lvl, val
838 else if (typ .eq. 6)
then
839 write(log_buf,
'(A15,I4,F12.6)')
'TAMG -POST', lvl, val
841 write(log_buf,
'(A15,I4,F12.6)')
'RESID', lvl, val
843 call neko_log%message(log_buf)
847 integer,
intent(in) :: nlvls
848 integer,
intent(in) :: smoo_type
851 character(len=LOG_SIZE) :: log_buf, smoo_name
853 call neko_log%section(
'PHMG')
855 if (smoo_type .eq. 1)
then
856 write(smoo_name,
'(A16)')
'CHEBY-acc JACOBI'
857 else if (smoo_type .eq. 2)
then
858 write(smoo_name,
'(A17)')
'CHEBY-acc SCHWARZ'
860 write(smoo_name,
'(A5)')
'CHEBY'
863 write(log_buf,
'(A28,I2,A8)') &
864 'Creating PHMG hierarchy with', &
866 call neko_log%message(log_buf)
870 write(log_buf,
'(A8,I2,A8,I2)') &
871 '-- level', i,
'-- lx:',
phmg%lvl(i)%Xh%lx
872 call neko_log%message(log_buf)
874 if (i .eq. clvl)
then
875 write(log_buf,
'(A19,A20)') &
877 call neko_log%message(log_buf)
879 write(log_buf,
'(A22,A20)') &
882 call neko_log%message(log_buf)
884 write(log_buf,
'(A28,I2)') &
886 phmg%lvl(i)%smoother_itrs
887 call neko_log%message(log_buf)
891 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.
integer, parameter, public coef_full
Retain every coefficient. The default, and the only scope that supports recompute_metrics(),...
integer, parameter, public coef_operator
Retain only what applying a discrete operator needs: , h1, h2, B and mult.
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_update_smoother_acc(this)
Rebuild the accelerator the Chebyshev smoother uses, which depends on the geometry.
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(this, coef, bclst, phmg_params)
subroutine phmg_update(this)
Bring the preconditioner back in sync after the mesh has changed.
subroutine phmg_update_smoother_eigs(this)
Mark every smoother to re-estimate its eigenvalues on the next solve.
subroutine phmg_update_coarse_geometry(this)
Sample the new fine coordinates on the coarse levels and rebuild their metrics. Level 0 shares the ca...
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 phmg_init_from_components(this, coef, bclst, smoother_itrs, cheby_acc, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree, pcrs_sched, update_enabled)
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.
Implements scalar_projector_t.
Overlapping schwarz solves.
Defines a function space.
integer, parameter, public gll
Implements multigrid using the TreeAMG hierarchy structure. USE:
subroutine, public neko_warning(warning_msg)
Reports a warning to standard output.
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.
Projector for scalar boundary conditions.
The function space for the SEM solution fields.
Type for the TreeAMG solver.