Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
pc_hsmg.f90
Go to the documentation of this file.
1! Copyright (c) 2008-2020, UCHICAGO ARGONNE, LLC.
2!
3! The UChicago Argonne, LLC as Operator of Argonne National
4! Laboratory holds copyright in the Software. The copyright holder
5! reserves all rights except those expressly granted to licensees,
6! and U.S. Government license rights.
7!
8! Redistribution and use in source and binary forms, with or without
9! modification, are permitted provided that the following conditions
10! are met:
11!
12! 1. Redistributions of source code must retain the above copyright
13! notice, this list of conditions and the disclaimer below.
14!
15! 2. Redistributions in binary form must reproduce the above copyright
16! notice, this list of conditions and the disclaimer (as noted below)
17! in the documentation and/or other materials provided with the
18! distribution.
19!
20! 3. Neither the name of ANL nor the names of its contributors
21! may be used to endorse or promote products derived from this software
22! without specific prior written permission.
23!
24! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
25! "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
26! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
27! FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL
28! UCHICAGO ARGONNE, LLC, THE U.S. DEPARTMENT OF
29! ENERGY OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
30! SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED
31! TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
32! DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
33! THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
34! (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
35! OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
36!
37! Additional BSD Notice
38! ---------------------
39! 1. This notice is required to be provided under our contract with
40! the U.S. Department of Energy (DOE). This work was produced at
41! Argonne National Laboratory under Contract
42! No. DE-AC02-06CH11357 with the DOE.
43!
44! 2. Neither the United States Government nor UCHICAGO ARGONNE,
45! LLC nor any of their employees, makes any warranty,
46! express or implied, or assumes any liability or responsibility for the
47! accuracy, completeness, or usefulness of any information, apparatus,
48! product, or process disclosed, or represents that its use would not
49! infringe privately-owned rights.
50!
51! 3. Also, reference herein to any specific commercial products, process,
52! or services by trade name, trademark, manufacturer or otherwise does
53! not necessarily constitute or imply its endorsement, recommendation,
54! or favoring by the United States Government or UCHICAGO ARGONNE LLC.
55! The views and opinions of authors expressed
56! herein do not necessarily state or reflect those of the United States
57! Government or UCHICAGO ARGONNE, LLC, and shall
58! not be used for advertising or product endorsement purposes.
59!
61module hsmg
63 use num_types, only : rp
64 use math, only : copy, col2, add2
65 use utils, only : neko_error
66 use precon, only : pc_t, precon_allocator, precon_destroy
67 use ax_product, only : ax_t, ax_helm_allocator
68 use gather_scatter, only : gs_t, gs_op_add
70 use bc, only : bc_t
71 use bc_list, only : bc_list_t
73 use dirichlet, only : dirichlet_t
74 use schwarz, only : schwarz_t
75 use jacobi, only : jacobi_t
76 use sx_jacobi, only : sx_jacobi_t
82 use space, only : space_t, gll
83 use dofmap, only : dofmap_t
84 use field, only : field_t
85 use coefs, only : coef_t, coef_operator
86 use mesh, only : mesh_t
87 use json_module, only : json_file
89 use krylov, only : ksp_t, ksp_monitor_t, ksp_max_iter, &
90 krylov_solver_factory
93 use logger, only : neko_log, log_size
94 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
95 !$ use omp_lib
96 implicit none
97 private
98
99 !Struct to arrange our multigridlevels
100 type, private :: multigrid_t
101 type(dofmap_t), pointer :: dof => null()
102 type(gs_t), pointer :: gs_h => null()
103 type(space_t), pointer :: xh => null()
104 type(coef_t), pointer :: coef => null()
105 type(scalar_bc_projector_t), pointer :: bc_projector => null()
106 type(schwarz_t), pointer :: schwarz => null()
107 type(field_t), pointer :: e => null()
108 end type multigrid_t
109
110 type, public, extends(pc_t) :: hsmg_t
111 type(mesh_t), pointer :: msh => null()
112 integer :: nlvls
113 type(multigrid_t), allocatable :: grids(:)
114 type(gs_t) :: gs_crs, gs_mg
115 type(space_t) :: xh_crs, xh_mg
116 type(dofmap_t) :: dm_crs, dm_mg
117 type(coef_t) :: c_crs, c_mg
118 type(zero_dirichlet_t) :: bc_crs, bc_mg, bc_reg
119 type(scalar_bc_projector_t) :: bc_projector_crs, bc_projector_mg, &
120 bc_projector_reg
121 type(schwarz_t) :: schwarz, schwarz_mg, schwarz_crs
123 type(field_t) :: e, e_mg, e_crs
124 type(field_t) :: wf
125 class(ksp_t), allocatable :: crs_solver
126 type(tamg_solver_t), allocatable :: amg_solver
127 integer :: niter
128 class(pc_t), allocatable :: pc_crs
129 class(ax_t), allocatable :: ax
130 real(kind=rp), allocatable :: r(:)
131 type(interpolator_t) :: interp_fine_mid
132 type(interpolator_t) :: interp_mid_crs
133 real(kind=rp), allocatable :: w(:)
134 type(c_ptr) :: w_d = c_null_ptr
135 type(c_ptr) :: r_d = c_null_ptr
136 type(c_ptr) :: hsmg_event = c_null_ptr
137 type(c_ptr) :: gs_event = c_null_ptr
138 contains
139 procedure, pass(this) :: init => hsmg_init
140 procedure, pass(this) :: init_from_components => &
142 procedure, pass(this) :: free => hsmg_free
143 procedure, pass(this) :: solve => hsmg_solve
144 procedure, pass(this) :: update => hsmg_set_h
145 end type hsmg_t
146
147contains
148
149 subroutine hsmg_init(this, coef, bclst, hsmg_params)
150 class(hsmg_t), intent(inout), target :: this
151 type(coef_t), intent(in), target :: coef
152 type(bc_list_t), intent(inout), target :: bclst
153 type(json_file), intent(inout) :: hsmg_params
154 character(len=:), allocatable :: crs_solver, crs_pc
155 logical :: crs_monitor
156 integer :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
157
158 ! Exract coarse grid parameters
159
160 ! Common parameters for the coarse grid
161 call json_get_or_default(hsmg_params, 'coarse_grid.solver', &
162 crs_solver, "cg")
163
164 !
165 ! Parameters for a Krylov based coarse grid solverthis
166 !
167 call json_get_or_default(hsmg_params, 'coarse_grid.iterations', &
168 this%niter, 10)
169
170 call json_get_or_default(hsmg_params, 'coarse_grid.preconditioner', &
171 crs_pc, "jacobi")
172
173 call json_get_or_default(hsmg_params, 'coarse_grid.monitor', &
174 crs_monitor, .false.)
175
176 !
177 ! Parameters for a tree-amg based coarse grid solver
178 !
179 call json_get_or_default(hsmg_params, 'coarse_grid.levels', &
180 crs_tamg_lvls, 3)
181
182 call json_get_or_default(hsmg_params, 'coarse_grid.iterations', &
183 crs_tamg_itrs, 1)
184
185 call json_get_or_default(hsmg_params, 'coarse_grid.cheby_degree', &
186 crs_tamg_cheby_degree, 4)
187
188 call this%init_from_components(coef, bclst, crs_solver, crs_pc, &
189 crs_monitor, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
190
191 end subroutine hsmg_init
192
193 subroutine hsmg_init_from_components(this, coef, bclst, crs_solver, crs_pc, &
194 crs_monitor, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
195 class(hsmg_t), intent(inout), target :: this
196 type(coef_t), intent(in), target :: coef
197 type(bc_list_t), intent(inout), target :: bclst
198 character(len=:), intent(inout), allocatable :: crs_solver, crs_pc
199 logical, intent(inout) :: crs_monitor
200 integer, intent(in) :: crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree
201 integer :: n, i
202 integer :: lx_crs, lx_mid
203 class(bc_t), pointer :: bc_i
204
205 character(len=LOG_SIZE) :: log_buf
206
207 call this%free()
210 this%nlvls = 3
211 lx_crs = 2
212 if (coef%Xh%lx .lt. 5) then
213 lx_mid = max(coef%Xh%lx-1,3)
214
215 if (coef%Xh%lx .le. 2) then
216 call neko_error('Polynomial order < 2 not supported for hsmg precon')
217 end if
218
219 else
220 lx_mid = 4
221 end if
222
223
224 call neko_log%section('HSMG')
225 if (this%nlvls .lt. 1e1) then
226 write(log_buf, '(A,I1,A)') 'HSMG hierarchy : ', &
227 this%nlvls, ' levels'
228 else if (this%nlvls .lt. 1e2) then
229 write(log_buf, '(A,I2,A)') 'HSMG hierarchy : ', &
230 this%nlvls, ' levels'
231 else if (this%nlvls .lt. 1e3) then
232 write(log_buf, '(A,I3,A)') 'HSMG hierarchy : ', this%nlvls, &
233 ' levels'
234 else
235 write(log_buf, '(A,I6,A)') 'HSMG hierarchy : ', this%nlvls, &
236 ' levels'
237 end if
238 call neko_log%message(log_buf)
239 if (trim(crs_solver) .ne. 'tamg' .or. trim(crs_solver) .eq. 'cheby') then
240 call neko_log%message('Coarse grid solver : (' // trim(crs_solver) // &
241 ', ' // trim(crs_pc) // ')')
242
243 if (this%niter .lt. 1e1) then
244 write(log_buf, '(A,I1)') 'Coarse grid iters. : ', this%niter
245 else if (this%niter .lt. 1e2) then
246 write(log_buf, '(A,I2)') 'Coarse grid iters. : ', this%niter
247 else if (this%niter .lt. 1e3) then
248 write(log_buf, '(A,I3)') 'Coarse grid iters. : ', this%niter
249 else if (this%niter .lt. 1e4) then
250 write(log_buf, '(A,I4)') 'Coarse grid iters. : ', this%niter
251 else
252 write(log_buf, '(A,I6)') 'Coarse grid iters. : ', this%niter
253 end if
254
255 call neko_log%message(log_buf)
256 else
257 call neko_log%message('Coarse grid solver : ' // trim(crs_solver) )
258 end if
259
260 this%msh => coef%msh
261 allocate(this%grids(this%nlvls))
262 allocate(this%w(coef%dof%size()))
263 allocate(this%r(coef%dof%size()))
264
265
266 ! Compute all elements as if they are deformed
267 call coef%msh%all_deformed()
268
269 n = coef%dof%size()
270 call this%e%init(coef%dof, 'work array')
271 call this%wf%init(coef%dof, 'work 2')
272
273 call this%Xh_crs%init(gll, lx_crs, lx_crs, lx_crs)
274 call this%dm_crs%init(coef%msh, this%Xh_crs)
275 call this%gs_crs%init(this%dm_crs)
276 call this%e_crs%init(this%dm_crs, 'work crs')
277 ! Both coarse levels read G_ij, h1, h2, B and mult and nothing else, and
278 ! hsmg never rebuilds their geometry, so the rest is dead weight.
279 call this%c_crs%init(this%gs_crs, coef_operator)
280
281 call this%Xh_mg%init(gll, lx_mid, lx_mid, lx_mid)
282 call this%dm_mg%init(coef%msh, this%Xh_mg)
283 call this%gs_mg%init(this%dm_mg)
284 call this%e_mg%init(this%dm_mg, 'work midl')
285 call this%c_mg%init(this%gs_mg, coef_operator)
286
287 ! Create backend specific Ax operator
288 call ax_helm_allocator(this%ax, type_name = "standard")
289
290 call this%bc_crs%init_base(this%c_crs)
291 call this%bc_mg%init_base(this%c_mg)
292 call this%bc_reg%init_base(coef)
293 if (bclst%size() .gt. 0) then
294 do i = 1, bclst%size()
295 bc_i => bclst%get(i)
296 call this%bc_reg%mark_facets(bc_i%marked_facet)
297 bc_i => bclst%get(i)
298 call this%bc_crs%mark_facets(bc_i%marked_facet)
299 bc_i => bclst%get(i)
300 call this%bc_mg%mark_facets(bc_i%marked_facet)
301 end do
302 end if
303 call this%bc_reg%finalize()
304 call this%bc_crs%finalize()
305 call this%bc_mg%finalize()
306
307 call this%bc_projector_reg%mark(this%bc_reg)
308 call this%bc_projector_crs%mark(this%bc_crs)
309 call this%bc_projector_mg%mark(this%bc_mg)
310
311 call this%schwarz%init(coef%Xh, coef%dof, coef%gs_h, &
312 this%bc_projector_reg, coef%msh)
313 call this%schwarz_mg%init(this%Xh_mg, this%dm_mg, this%gs_mg,&
314 this%bc_projector_mg, coef%msh)
315
316 call this%interp_fine_mid%init(coef%Xh, this%Xh_mg)
317 call this%interp_mid_crs%init(this%Xh_mg, this%Xh_crs)
318
319 call hsmg_fill_grid(coef%dof, coef%gs_h, coef%Xh, coef, &
320 this%bc_projector_reg, this%schwarz, this%e, this%grids, 3)
321 call hsmg_fill_grid(this%dm_mg, this%gs_mg, this%Xh_mg, this%c_mg, &
322 this%bc_projector_mg, this%schwarz_mg, this%e_mg, this%grids, 2)
323 call hsmg_fill_grid(this%dm_crs, this%gs_crs, this%Xh_crs, &
324 this%c_crs, this%bc_projector_crs, this%schwarz_crs, &
325 this%e_crs, this%grids, 1)
326
327 call hsmg_set_h(this)
328 if (neko_bcknd_device .eq. 1) then
329 call device_map(this%w, this%w_d, n)
330 call device_map(this%r, this%r_d, n)
331 end if
332
333 call device_event_create(this%hsmg_event, 2)
334 call device_event_create(this%gs_event, 2)
335
336
337
338 ! Create a backend specific krylov solver
339 if (trim(crs_solver) .eq. 'tamg') then
340 allocate(this%amg_solver)
341 call this%amg_solver%init(this%ax, this%grids(1)%e%Xh, &
342 this%grids(1)%coef, this%msh, this%grids(1)%gs_h, crs_tamg_lvls, &
343 this%grids(1)%bc_projector, crs_tamg_itrs, crs_tamg_cheby_degree)
344 else
345 ! Create a backend specific preconditioner
346 call precon_allocator(this%pc_crs, crs_pc)
347
348 select type (pc => this%pc_crs)
349 type is (jacobi_t)
350 call pc%init(this%c_crs, this%dm_crs, this%gs_crs)
351 type is (sx_jacobi_t)
352 call pc%init(this%c_crs, this%dm_crs, this%gs_crs)
353 type is (device_jacobi_t)
354 call pc%init(this%c_crs, this%dm_crs, this%gs_crs)
355 end select
356
357 call krylov_solver_factory(this%crs_solver, &
358 this%dm_crs%size(), trim(crs_solver), ksp_max_iter, &
359 m = this%pc_crs, monitor = crs_monitor)
360 end if
361
362 call neko_log%end_section()
363
364 end subroutine hsmg_init_from_components
365
366 subroutine hsmg_set_h(this)
367 class(hsmg_t), intent(inout) :: this
368 ! integer :: i
369 ! Yeah I dont really know what to do here. For incompressible flow not
370 ! much happens
371 this%grids(1)%coef%ifh2 = .false.
372 call copy(this%grids(1)%coef%h1, this%grids(3)%coef%h1, &
373 this%grids(1)%dof%size())
374 if (neko_bcknd_device .eq. 1) then
375 call device_copy(this%grids(1)%coef%h1_d, this%grids(3)%coef%h1_d, &
376 this%grids(1)%dof%size())
377 end if
378 end subroutine hsmg_set_h
379
380
381 subroutine hsmg_fill_grid(dof, gs_h, Xh, coef, bc_projector, schwarz, e, grids, l)
382 type(dofmap_t), target, intent(in) :: dof
383 type(gs_t), target, intent(in) :: gs_h
384 type(space_t), target, intent(in) :: Xh
385 type(coef_t), target, intent(in) :: coef
386 type(scalar_bc_projector_t), target, intent(in) :: bc_projector
387 type(schwarz_t), target, intent(in) :: schwarz
388 type(field_t), target, intent(in) :: e
389 integer, intent(in) :: l
390 type(multigrid_t), intent(inout), dimension(l) :: grids
391
392
393 grids(l)%dof => dof
394 grids(l)%gs_h => gs_h
395 grids(l)%Xh => xh
396 grids(l)%coef => coef
397 grids(l)%bc_projector => bc_projector
398 grids(l)%schwarz => schwarz
399 grids(l)%e => e
400
401 end subroutine hsmg_fill_grid
402
403 subroutine hsmg_free(this)
404 class(hsmg_t), intent(inout) :: this
405
406 if (allocated(this%ax)) then
407 deallocate(this%ax)
408 end if
409
410 if (allocated(this%grids)) then
411 deallocate(this%grids)
412 end if
413
414 if (allocated(this%w)) then
415 if (c_associated(this%w_d)) then
416 call device_unmap(this%w, this%w_d)
417 end if
418 deallocate(this%w)
419 end if
420
421 if (allocated(this%r)) then
422 if (c_associated(this%r_d)) then
423 call device_unmap(this%r, this%r_d)
424 end if
425 deallocate(this%r)
426 end if
427
428 call this%schwarz%free()
429 call this%schwarz_mg%free()
430
431 call this%c_crs%free()
432 call this%c_mg%free()
433 call this%e%free()
434 call this%e_mg%free()
435 call this%e_crs%free()
436 call this%wf%free()
437
438 call this%gs_crs%free()
439 call this%gs_mg%free()
440 call this%interp_mid_crs%free()
441 call this%interp_fine_mid%free()
442
443 call this%bc_crs%free()
444 call this%bc_mg%free()
445 call this%bc_reg%free()
446
447 call this%bc_projector_reg%free()
448 call this%bc_projector_crs%free()
449 call this%bc_projector_mg%free()
450
451 if (allocated(this%crs_solver)) then
452 call this%crs_solver%free()
453 deallocate(this%crs_solver)
454 end if
455
456 if (allocated(this%amg_solver)) then
457 call this%amg_solver%free()
458 deallocate(this%amg_solver)
459 end if
460
461 if (allocated(this%pc_crs)) then
462 call precon_destroy(this%pc_crs)
463 end if
464
465 if (c_associated(this%hsmg_event)) then
466 call device_event_destroy(this%hsmg_event)
467 end if
468 if (c_associated(this%gs_event)) then
469 call device_event_destroy(this%gs_event)
470 end if
471
472 call this%dm_crs%free()
473 call this%dm_mg%free()
474 call this%Xh_crs%free()
475 call this%Xh_mg%free()
476
477 end subroutine hsmg_free
478
480 subroutine hsmg_solve(this, z, r, n)
481 integer, intent(in) :: n
482 class(hsmg_t), intent(inout) :: this
483 real(kind=rp), dimension(n), intent(inout) :: z
484 real(kind=rp), dimension(n), intent(inout) :: r
485 type(c_ptr) :: z_d, r_d
486 type(ksp_monitor_t) :: crs_info
487 integer :: thrdid, nthrds
488
489 call profiler_start_region('HSMG_solve', 8)
490 if (neko_bcknd_device .eq. 1) then
491 z_d = device_get_ptr(z)
492 r_d = device_get_ptr(r)
493 !We should not work with the input
494 call device_copy(this%r_d, r_d, n)
495 call this%bc_projector_reg%apply(this%r, n)
496
497 !OVERLAPPING Schwarz exchange and solve
498 !! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
499 call device_col2(this%r_d, this%grids(3)%coef%mult_d, &
500 this%grids(3)%dof%size())
501 !Restrict to middle level
502 call this%interp_fine_mid%map(this%e%x, this%r, &
503 this%msh%nelv, this%grids(2)%Xh)
504 call this%grids(2)%gs_h%op(this%e%x, &
505 this%grids(2)%dof%size(), gs_op_add, this%gs_event)
506 call device_event_sync(this%gs_event)
507 !This should probably be double checked again
508 call device_copy(this%r_d, r_d, n)
509 call this%bc_projector_reg%apply(this%r, n)
510 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
511 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
512 !OVERLAPPING Schwarz exchange and solve
513 call device_col2(this%w_d, this%grids(2)%coef%mult_d, &
514 this%grids(2)%dof%size())
515 !restrict residual to crs
516 call this%interp_mid_crs%map(this%wf%x, this%w, this%msh%nelv, &
517 this%grids(1)%Xh)
518 !Crs solve
519 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
520 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
521
522 !$omp parallel private(thrdid, nthrds)
523
524 thrdid = 0
525 nthrds = 1
526 !$ thrdid = omp_get_thread_num()
527 !$ nthrds = omp_get_num_threads()
528
529 if (thrdid .eq. 0) then
530 call profiler_start_region('HSMG_schwarz', 9)
531 call this%grids(3)%schwarz%compute(z, this%r)
532 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
533 call profiler_end_region('HSMG_schwarz', 9)
534 end if
535 if (nthrds .eq. 1 .or. thrdid .eq. 1) then
536 call profiler_start_region('HSMG_coarse_grid', 10)
537 call this%grids(1)%gs_h%op(this%wf%x, &
538 this%grids(1)%dof%size(), gs_op_add, this%gs_event)
539 call device_event_sync(this%gs_event)
540 call this%grids(1)%bc_projector%apply(this%wf%x, &
541 this%grids(1)%dof%size())
542 call profiler_start_region('HSMG_coarse_solve', 11)
543 if (allocated(this%amg_solver)) then
544 call this%amg_solver%solve(this%grids(1)%e%x, this%wf%x, &
545 this%grids(1)%dof%size())
546 else
547 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, &
548 this%wf%x, &
549 this%grids(1)%dof%size(), &
550 this%grids(1)%coef, &
551 this%grids(1)%bc_projector, &
552 this%grids(1)%gs_h, this%niter)
553 end if
554 call profiler_end_region('HSMG_coarse_solve', 11)
555 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x,&
556 this%grids(1)%dof%size())
557 call profiler_end_region('HSMG_coarse_grid', 10)
558 end if
559 !$omp end parallel
560
561 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
562 this%msh%nelv, this%grids(2)%Xh)
563 call device_add2(this%grids(2)%e%x_d, this%w_d, this%grids(2)%dof%size())
564
565 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
566 this%msh%nelv, this%grids(3)%Xh)
567 call device_add2(z_d, this%w_d, this%grids(3)%dof%size())
568 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), &
569 gs_op_add, this%gs_event)
570 call device_event_sync(this%gs_event)
571 call device_col2(z_d, this%grids(3)%coef%mult_d, &
572 this%grids(3)%dof%size())
573 else
574 !We should not work with the input
575 call copy(this%r, r, n)
576
577 !OVERLAPPING Schwarz exchange and solve
578 call this%grids(3)%schwarz%compute(z, this%r)
579 ! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
580 call col2(this%r, this%grids(3)%coef%mult, &
581 this%grids(3)%dof%size())
582 !Restrict to middle level
583 call this%interp_fine_mid%map(this%w, this%r, &
584 this%msh%nelv, this%grids(2)%Xh)
585 call this%grids(2)%gs_h%op(this%w, this%grids(2)%dof%size(), gs_op_add)
586 !OVERLAPPING Schwarz exchange and solve
587 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
588 call col2(this%w, this%grids(2)%coef%mult, this%grids(2)%dof%size())
589 !restrict residual to crs
590 call this%interp_mid_crs%map(this%r, this%w, &
591 this%msh%nelv, this%grids(1)%Xh)
592 !Crs solve
593
594 call this%grids(1)%gs_h%op(this%r, this%grids(1)%dof%size(), gs_op_add)
595 call this%grids(1)%bc_projector%apply(this%r, this%grids(1)%dof%size())
596
597 call profiler_start_region('HSMG_coarse-solve', 11)
598 if (allocated(this%amg_solver)) then
599 call this%amg_solver%solve(this%grids(1)%e%x, this%r, &
600 this%grids(1)%dof%size())
601 else
602 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, this%r, &
603 this%grids(1)%dof%size(), &
604 this%grids(1)%coef, &
605 this%grids(1)%bc_projector, &
606 this%grids(1)%gs_h, this%niter)
607 end if
608 call profiler_end_region('HSMG_coarse-solve', 11)
609
610 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x, &
611 this%grids(1)%dof%size())
612
613
614 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
615 this%msh%nelv, this%grids(2)%Xh)
616 call add2(this%grids(2)%e%x, this%w, this%grids(2)%dof%size())
617
618 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
619 this%msh%nelv, this%grids(3)%Xh)
620 call add2(z, this%w, this%grids(3)%dof%size())
621 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), gs_op_add)
622 call col2(z, this%grids(3)%coef%mult, this%grids(3)%dof%size())
623
624 end if
625 call profiler_end_region('HSMG_solve', 8)
626 end subroutine hsmg_solve
627end module hsmg
__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.
Definition device.F90:113
Map a Fortran array to a device (allocate and associate)
Definition device.F90:83
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Retrieves a parameter by name or assigns a provided default value. In the latter case also adds the m...
Defines a Matrix-vector product.
Definition ax.f90:34
Defines a list of bc_t.
Definition bc_list.f90:34
Defines a boundary condition.
Definition bc.f90:34
Coefficients.
Definition coef.f90:34
integer, parameter, public coef_operator
Retain only what applying a discrete operator needs: , h1, h2, B and mult.
Definition coef.f90:89
Jacobi preconditioner accelerator backend.
subroutine, public device_add2(a_d, b_d, n, strm)
Vector addition .
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 .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
subroutine, public device_event_sync(event)
Synchronize an event.
Definition device.F90:1667
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1623
subroutine, public device_event_create(event, flags)
Create a device event queue.
Definition device.F90:1589
Defines a dirichlet boundary condition.
Definition dirichlet.f90:34
Defines a mapping of the degrees of freedom.
Definition dofmap.f90:35
Defines a field.
Definition field.f90:34
Gather-scatter.
Krylov preconditioner.
Definition pc_hsmg.f90:61
subroutine hsmg_init(this, coef, bclst, hsmg_params)
Definition pc_hsmg.f90:150
subroutine hsmg_solve(this, z, r, n)
The h1mg preconditioner from Nek5000.
Definition pc_hsmg.f90:481
subroutine hsmg_set_h(this)
Definition pc_hsmg.f90:367
subroutine hsmg_free(this)
Definition pc_hsmg.f90:404
subroutine hsmg_init_from_components(this, coef, bclst, crs_solver, crs_pc, crs_monitor, crs_tamg_lvls, crs_tamg_itrs, crs_tamg_cheby_degree)
Definition pc_hsmg.f90:195
subroutine hsmg_fill_grid(dof, gs_h, xh, coef, bc_projector, schwarz, e, grids, l)
Definition pc_hsmg.f90:382
Routines to interpolate between different spaces.
Jacobi preconditioner.
Definition pc_jacobi.f90:34
Utilities for retrieving parameters from the case files.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
integer, parameter, public ksp_max_iter
Maximum number of iters.
Definition krylov.f90:52
Logging routines.
Definition log.f90:34
type(log_t), public neko_log
Global log stream.
Definition log.f90:80
integer, parameter, public log_size
Definition log.f90:46
Definition math.f90:60
subroutine, public add2(a, b, n)
Vector addition .
Definition math.f90:903
subroutine, public col2(a, b, n)
Vector multiplication .
Definition math.f90:1049
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:294
Defines a mesh.
Definition mesh.f90:34
Build configurations.
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Krylov preconditioner.
Definition precon.f90:34
Profiling interface.
Definition profiler.F90:34
subroutine, public profiler_start_region(name, region_id)
Started a named (name) profiler region.
Definition profiler.F90:79
subroutine, public profiler_end_region(name, region_id)
End the most recently started profiler region.
Definition profiler.F90:116
Implements scalar_projector_t.
Overlapping schwarz solves.
Definition schwarz.f90:61
Defines a function space.
Definition space.f90:34
integer, parameter, public gll
Definition space.f90:50
Jacobi preconditioner SX-Aurora backend.
Implements multigrid using the TreeAMG hierarchy structure. USE:
Utilities.
Definition utils.f90:35
Defines a zero-valued Dirichlet boundary condition.
Base type for a matrix-vector product providing .
Definition ax.f90:43
Base type for a boundary condition.
Definition bc.f90:72
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:93
Defines a jacobi preconditioner.
Generic Dirichlet boundary condition on .
Definition dirichlet.f90:49
Gather-scatter kernel.
Interpolation between two space::space_t.
Defines a jacobi preconditioner.
Definition pc_jacobi.f90:45
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:57
Base abstract type for a canonical Krylov method, solving .
Definition krylov.f90:74
Defines a canonical Krylov preconditioner.
Definition precon.f90:40
Projector for scalar boundary conditions.
The function space for the SEM solution fields.
Definition space.f90:64
Defines a jacobi preconditioner for SX-Aurora.
Type for the TreeAMG solver.
Zero-valued Dirichlet boundary condition. Used for no-slip walls, but also for various auxillary cond...
#define max(a, b)
Definition tensor.cu:40