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 call this%ax%free()
408 deallocate(this%ax)
409 end if
410
411 if (allocated(this%grids)) then
412 deallocate(this%grids)
413 end if
414
415 if (allocated(this%w)) then
416 if (c_associated(this%w_d)) then
417 call device_unmap(this%w, this%w_d)
418 end if
419 deallocate(this%w)
420 end if
421
422 if (allocated(this%r)) then
423 if (c_associated(this%r_d)) then
424 call device_unmap(this%r, this%r_d)
425 end if
426 deallocate(this%r)
427 end if
428
429 call this%schwarz%free()
430 call this%schwarz_mg%free()
431
432 call this%c_crs%free()
433 call this%c_mg%free()
434 call this%e%free()
435 call this%e_mg%free()
436 call this%e_crs%free()
437 call this%wf%free()
438
439 call this%gs_crs%free()
440 call this%gs_mg%free()
441 call this%interp_mid_crs%free()
442 call this%interp_fine_mid%free()
443
444 call this%bc_crs%free()
445 call this%bc_mg%free()
446 call this%bc_reg%free()
447
448 call this%bc_projector_reg%free()
449 call this%bc_projector_crs%free()
450 call this%bc_projector_mg%free()
451
452 if (allocated(this%crs_solver)) then
453 call this%crs_solver%free()
454 deallocate(this%crs_solver)
455 end if
456
457 if (allocated(this%amg_solver)) then
458 call this%amg_solver%free()
459 deallocate(this%amg_solver)
460 end if
461
462 if (allocated(this%pc_crs)) then
463 call precon_destroy(this%pc_crs)
464 end if
465
466 if (c_associated(this%hsmg_event)) then
467 call device_event_destroy(this%hsmg_event)
468 end if
469 if (c_associated(this%gs_event)) then
470 call device_event_destroy(this%gs_event)
471 end if
472
473 call this%dm_crs%free()
474 call this%dm_mg%free()
475 call this%Xh_crs%free()
476 call this%Xh_mg%free()
477
478 end subroutine hsmg_free
479
481 subroutine hsmg_solve(this, z, r, n)
482 integer, intent(in) :: n
483 class(hsmg_t), intent(inout) :: this
484 real(kind=rp), dimension(n), intent(inout) :: z
485 real(kind=rp), dimension(n), intent(inout) :: r
486 type(c_ptr) :: z_d, r_d
487 type(ksp_monitor_t) :: crs_info
488 integer :: thrdid, nthrds
489
490 call profiler_start_region('HSMG_solve', 8)
491 if (neko_bcknd_device .eq. 1) then
492 z_d = device_get_ptr(z)
493 r_d = device_get_ptr(r)
494 !We should not work with the input
495 call device_copy(this%r_d, r_d, n)
496 call this%bc_projector_reg%apply(this%r, n)
497
498 !OVERLAPPING Schwarz exchange and solve
499 !! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
500 call device_col2(this%r_d, this%grids(3)%coef%mult_d, &
501 this%grids(3)%dof%size())
502 !Restrict to middle level
503 call this%interp_fine_mid%map(this%e%x, this%r, &
504 this%msh%nelv, this%grids(2)%Xh)
505 call this%grids(2)%gs_h%op(this%e%x, &
506 this%grids(2)%dof%size(), gs_op_add, this%gs_event)
507 call device_event_sync(this%gs_event)
508 !This should probably be double checked again
509 call device_copy(this%r_d, r_d, n)
510 call this%bc_projector_reg%apply(this%r, n)
511 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
512 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
513 !OVERLAPPING Schwarz exchange and solve
514 call device_col2(this%w_d, this%grids(2)%coef%mult_d, &
515 this%grids(2)%dof%size())
516 !restrict residual to crs
517 call this%interp_mid_crs%map(this%wf%x, this%w, this%msh%nelv, &
518 this%grids(1)%Xh)
519 !Crs solve
520 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
521 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
522
523 !$omp parallel private(thrdid, nthrds)
524
525 thrdid = 0
526 nthrds = 1
527 !$ thrdid = omp_get_thread_num()
528 !$ nthrds = omp_get_num_threads()
529
530 if (thrdid .eq. 0) then
531 call profiler_start_region('HSMG_schwarz', 9)
532 call this%grids(3)%schwarz%compute(z, this%r)
533 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
534 call profiler_end_region('HSMG_schwarz', 9)
535 end if
536 if (nthrds .eq. 1 .or. thrdid .eq. 1) then
537 call profiler_start_region('HSMG_coarse_grid', 10)
538 call this%grids(1)%gs_h%op(this%wf%x, &
539 this%grids(1)%dof%size(), gs_op_add, this%gs_event)
540 call device_event_sync(this%gs_event)
541 call this%grids(1)%bc_projector%apply(this%wf%x, &
542 this%grids(1)%dof%size())
543 call profiler_start_region('HSMG_coarse_solve', 11)
544 if (allocated(this%amg_solver)) then
545 call this%amg_solver%solve(this%grids(1)%e%x, this%wf%x, &
546 this%grids(1)%dof%size())
547 else
548 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, &
549 this%wf%x, &
550 this%grids(1)%dof%size(), &
551 this%grids(1)%coef, &
552 this%grids(1)%bc_projector, &
553 this%grids(1)%gs_h, this%niter)
554 end if
555 call profiler_end_region('HSMG_coarse_solve', 11)
556 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x,&
557 this%grids(1)%dof%size())
558 call profiler_end_region('HSMG_coarse_grid', 10)
559 end if
560 !$omp end parallel
561
562 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
563 this%msh%nelv, this%grids(2)%Xh)
564 call device_add2(this%grids(2)%e%x_d, this%w_d, this%grids(2)%dof%size())
565
566 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
567 this%msh%nelv, this%grids(3)%Xh)
568 call device_add2(z_d, this%w_d, this%grids(3)%dof%size())
569 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), &
570 gs_op_add, this%gs_event)
571 call device_event_sync(this%gs_event)
572 call device_col2(z_d, this%grids(3)%coef%mult_d, &
573 this%grids(3)%dof%size())
574 else
575 !We should not work with the input
576 call copy(this%r, r, n)
577
578 !OVERLAPPING Schwarz exchange and solve
579 call this%grids(3)%schwarz%compute(z, this%r)
580 ! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
581 call col2(this%r, this%grids(3)%coef%mult, &
582 this%grids(3)%dof%size())
583 !Restrict to middle level
584 call this%interp_fine_mid%map(this%w, this%r, &
585 this%msh%nelv, this%grids(2)%Xh)
586 call this%grids(2)%gs_h%op(this%w, this%grids(2)%dof%size(), gs_op_add)
587 !OVERLAPPING Schwarz exchange and solve
588 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
589 call col2(this%w, this%grids(2)%coef%mult, this%grids(2)%dof%size())
590 !restrict residual to crs
591 call this%interp_mid_crs%map(this%r, this%w, &
592 this%msh%nelv, this%grids(1)%Xh)
593 !Crs solve
594
595 call this%grids(1)%gs_h%op(this%r, this%grids(1)%dof%size(), gs_op_add)
596 call this%grids(1)%bc_projector%apply(this%r, this%grids(1)%dof%size())
597
598 call profiler_start_region('HSMG_coarse-solve', 11)
599 if (allocated(this%amg_solver)) then
600 call this%amg_solver%solve(this%grids(1)%e%x, this%r, &
601 this%grids(1)%dof%size())
602 else
603 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, this%r, &
604 this%grids(1)%dof%size(), &
605 this%grids(1)%coef, &
606 this%grids(1)%bc_projector, &
607 this%grids(1)%gs_h, this%niter)
608 end if
609 call profiler_end_region('HSMG_coarse-solve', 11)
610
611 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x, &
612 this%grids(1)%dof%size())
613
614
615 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
616 this%msh%nelv, this%grids(2)%Xh)
617 call add2(this%grids(2)%e%x, this%w, this%grids(2)%dof%size())
618
619 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
620 this%msh%nelv, this%grids(3)%Xh)
621 call add2(z, this%w, this%grids(3)%dof%size())
622 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), gs_op_add)
623 call col2(z, this%grids(3)%coef%mult, this%grids(3)%dof%size())
624
625 end if
626 call profiler_end_region('HSMG_solve', 8)
627 end subroutine hsmg_solve
628end 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:131
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:482
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:91
integer, parameter, public log_size
Definition log.f90:46
Definition math.f90:60
subroutine, public add2(a, b, n)
Vector addition .
Definition math.f90:939
subroutine, public col2(a, b, n)
Vector multiplication .
Definition math.f90:1085
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:295
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:73
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:135
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