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
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, json)
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) :: json
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(json, '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(json, 'coarse_grid.iterations', &
168 this%niter, 10)
169
170 call json_get_or_default(json, 'coarse_grid.preconditioner', &
171 crs_pc, "jacobi")
172
173 call json_get_or_default(json, '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(json, 'coarse_grid.levels', &
180 crs_tamg_lvls, 3)
181
182 call json_get_or_default(json, 'coarse_grid.iterations', &
183 crs_tamg_itrs, 1)
184
185 call json_get_or_default(json, '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. The coarse levels compress their
288 ! geometric factors when the fine level does
289 if (coef%geo_compression) then
290 call this%c_crs%enable_geo_compression(coef%geo_compression_tol)
291 call this%c_mg%enable_geo_compression(coef%geo_compression_tol)
292 call ax_helm_allocator(this%ax, type_name = "standard_compr")
293 else
294 call ax_helm_allocator(this%ax, type_name = "standard")
295 end if
296
297 call this%bc_crs%init_base(this%c_crs)
298 call this%bc_mg%init_base(this%c_mg)
299 call this%bc_reg%init_base(coef)
300 if (bclst%size() .gt. 0) then
301 do i = 1, bclst%size()
302 bc_i => bclst%get(i)
303 call this%bc_reg%mark_facets(bc_i%marked_facet)
304 bc_i => bclst%get(i)
305 call this%bc_crs%mark_facets(bc_i%marked_facet)
306 bc_i => bclst%get(i)
307 call this%bc_mg%mark_facets(bc_i%marked_facet)
308 end do
309 end if
310 call this%bc_reg%finalize()
311 call this%bc_crs%finalize()
312 call this%bc_mg%finalize()
313
314 call this%bc_projector_reg%mark(this%bc_reg)
315 call this%bc_projector_crs%mark(this%bc_crs)
316 call this%bc_projector_mg%mark(this%bc_mg)
317
318 call this%schwarz%init(coef%Xh, coef%dof, coef%gs_h, &
319 this%bc_projector_reg, coef%msh)
320 call this%schwarz_mg%init(this%Xh_mg, this%dm_mg, this%gs_mg,&
321 this%bc_projector_mg, coef%msh)
322
323 call this%interp_fine_mid%init(coef%Xh, this%Xh_mg)
324 call this%interp_mid_crs%init(this%Xh_mg, this%Xh_crs)
325
326 call hsmg_fill_grid(coef%dof, coef%gs_h, coef%Xh, coef, &
327 this%bc_projector_reg, this%schwarz, this%e, this%grids, 3)
328 call hsmg_fill_grid(this%dm_mg, this%gs_mg, this%Xh_mg, this%c_mg, &
329 this%bc_projector_mg, this%schwarz_mg, this%e_mg, this%grids, 2)
330 call hsmg_fill_grid(this%dm_crs, this%gs_crs, this%Xh_crs, &
331 this%c_crs, this%bc_projector_crs, this%schwarz_crs, &
332 this%e_crs, this%grids, 1)
333
334 call hsmg_set_h(this)
335 if (neko_bcknd_device .eq. 1) then
336 call device_map(this%w, this%w_d, n)
337 call device_map(this%r, this%r_d, n)
338 end if
339
340 call device_event_create(this%hsmg_event, 2)
341 call device_event_create(this%gs_event, 2)
342
343
344
345 ! Create a backend specific krylov solver
346 if (trim(crs_solver) .eq. 'tamg') then
347 allocate(this%amg_solver)
348 call this%amg_solver%init(this%ax, this%grids(1)%e%Xh, &
349 this%grids(1)%coef, this%msh, this%grids(1)%gs_h, crs_tamg_lvls, &
350 this%grids(1)%bc_projector, crs_tamg_itrs, crs_tamg_cheby_degree)
351 else
352 ! Create a backend specific preconditioner
353 call precon_allocator(this%pc_crs, crs_pc)
354
355 select type (pc => this%pc_crs)
356 type is (jacobi_t)
357 call pc%init_from_components(this%c_crs, this%dm_crs, this%gs_crs)
358 type is (sx_jacobi_t)
359 call pc%init_from_components(this%c_crs, this%dm_crs, this%gs_crs)
360 type is (device_jacobi_t)
361 call pc%init_from_components(this%c_crs, this%dm_crs, this%gs_crs)
362 end select
363
364 call krylov_solver_factory(this%crs_solver, &
365 this%dm_crs%size(), trim(crs_solver), ksp_max_iter, &
366 m = this%pc_crs, monitor = crs_monitor)
367 end if
368
369 call neko_log%end_section()
370
371 end subroutine hsmg_init_from_components
372
373 subroutine hsmg_set_h(this)
374 class(hsmg_t), intent(inout) :: this
375 ! integer :: i
376 ! Yeah I dont really know what to do here. For incompressible flow not
377 ! much happens
378 this%grids(1)%coef%ifh2 = .false.
379 call copy(this%grids(1)%coef%h1, this%grids(3)%coef%h1, &
380 this%grids(1)%dof%size())
381 if (neko_bcknd_device .eq. 1) then
382 call device_copy(this%grids(1)%coef%h1_d, this%grids(3)%coef%h1_d, &
383 this%grids(1)%dof%size())
384 end if
385 end subroutine hsmg_set_h
386
387
388 subroutine hsmg_fill_grid(dof, gs_h, Xh, coef, bc_projector, schwarz, e, grids, l)
389 type(dofmap_t), target, intent(in) :: dof
390 type(gs_t), target, intent(in) :: gs_h
391 type(space_t), target, intent(in) :: Xh
392 type(coef_t), target, intent(in) :: coef
393 type(scalar_bc_projector_t), target, intent(in) :: bc_projector
394 type(schwarz_t), target, intent(in) :: schwarz
395 type(field_t), target, intent(in) :: e
396 integer, intent(in) :: l
397 type(multigrid_t), intent(inout), dimension(l) :: grids
398
399
400 grids(l)%dof => dof
401 grids(l)%gs_h => gs_h
402 grids(l)%Xh => xh
403 grids(l)%coef => coef
404 grids(l)%bc_projector => bc_projector
405 grids(l)%schwarz => schwarz
406 grids(l)%e => e
407
408 end subroutine hsmg_fill_grid
409
410 subroutine hsmg_free(this)
411 class(hsmg_t), intent(inout) :: this
412
413 if (allocated(this%ax)) then
414 call this%ax%free()
415 deallocate(this%ax)
416 end if
417
418 if (allocated(this%grids)) then
419 deallocate(this%grids)
420 end if
421
422 if (allocated(this%w)) then
423 if (c_associated(this%w_d)) then
424 call device_unmap(this%w, this%w_d)
425 end if
426 deallocate(this%w)
427 end if
428
429 if (allocated(this%r)) then
430 if (c_associated(this%r_d)) then
431 call device_unmap(this%r, this%r_d)
432 end if
433 deallocate(this%r)
434 end if
435
436 call this%schwarz%free()
437 call this%schwarz_mg%free()
438
439 call this%c_crs%free()
440 call this%c_mg%free()
441 call this%e%free()
442 call this%e_mg%free()
443 call this%e_crs%free()
444 call this%wf%free()
445
446 call this%gs_crs%free()
447 call this%gs_mg%free()
448 call this%interp_mid_crs%free()
449 call this%interp_fine_mid%free()
450
451 call this%bc_crs%free()
452 call this%bc_mg%free()
453 call this%bc_reg%free()
454
455 call this%bc_projector_reg%free()
456 call this%bc_projector_crs%free()
457 call this%bc_projector_mg%free()
458
459 if (allocated(this%crs_solver)) then
460 call this%crs_solver%free()
461 deallocate(this%crs_solver)
462 end if
463
464 if (allocated(this%amg_solver)) then
465 call this%amg_solver%free()
466 deallocate(this%amg_solver)
467 end if
468
469 if (allocated(this%pc_crs)) then
470 call this%pc_crs%free()
471 end if
472
473 if (c_associated(this%hsmg_event)) then
474 call device_event_destroy(this%hsmg_event)
475 end if
476 if (c_associated(this%gs_event)) then
477 call device_event_destroy(this%gs_event)
478 end if
479
480 call this%dm_crs%free()
481 call this%dm_mg%free()
482 call this%Xh_crs%free()
483 call this%Xh_mg%free()
484
485 end subroutine hsmg_free
486
488 subroutine hsmg_solve(this, z, r, n)
489 integer, intent(in) :: n
490 class(hsmg_t), intent(inout) :: this
491 real(kind=rp), dimension(n), intent(inout) :: z
492 real(kind=rp), dimension(n), intent(inout) :: r
493 type(c_ptr) :: z_d, r_d
494 type(ksp_monitor_t) :: crs_info
495 integer :: thrdid, nthrds
496
497 call profiler_start_region('HSMG_solve', 8)
498 if (neko_bcknd_device .eq. 1) then
499 z_d = device_get_ptr(z)
500 r_d = device_get_ptr(r)
501 !We should not work with the input
502 call device_copy(this%r_d, r_d, n)
503 call this%bc_projector_reg%apply(this%r, n)
504
505 !OVERLAPPING Schwarz exchange and solve
506 !! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
507 call device_col2(this%r_d, this%grids(3)%coef%mult_d, &
508 this%grids(3)%dof%size())
509 !Restrict to middle level
510 call this%interp_fine_mid%map(this%e%x, this%r, &
511 this%msh%nelv, this%grids(2)%Xh)
512 call this%grids(2)%gs_h%op(this%e%x, &
513 this%grids(2)%dof%size(), gs_op_add, this%gs_event)
514 call device_event_sync(this%gs_event)
515 !This should probably be double checked again
516 call device_copy(this%r_d, r_d, n)
517 call this%bc_projector_reg%apply(this%r, n)
518 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
519 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
520 !OVERLAPPING Schwarz exchange and solve
521 call device_col2(this%w_d, this%grids(2)%coef%mult_d, &
522 this%grids(2)%dof%size())
523 !restrict residual to crs
524 call this%interp_mid_crs%map(this%wf%x, this%w, this%msh%nelv, &
525 this%grids(1)%Xh)
526 !Crs solve
527 call device_copy(this%w_d, this%e%x_d, this%grids(2)%dof%size())
528 call this%bc_projector_mg%apply(this%w, this%grids(2)%dof%size())
529
530 !$omp parallel private(thrdid, nthrds)
531
532 thrdid = 0
533 nthrds = 1
534 !$ thrdid = omp_get_thread_num()
535 !$ nthrds = omp_get_num_threads()
536
537 if (thrdid .eq. 0) then
538 call profiler_start_region('HSMG_schwarz', 9)
539 call this%grids(3)%schwarz%compute(z, this%r)
540 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
541 call profiler_end_region('HSMG_schwarz', 9)
542 end if
543 if (nthrds .eq. 1 .or. thrdid .eq. 1) then
544 call profiler_start_region('HSMG_coarse_grid', 10)
545 call this%grids(1)%gs_h%op(this%wf%x, &
546 this%grids(1)%dof%size(), gs_op_add, this%gs_event)
547 call device_event_sync(this%gs_event)
548 call this%grids(1)%bc_projector%apply(this%wf%x, &
549 this%grids(1)%dof%size())
550 call profiler_start_region('HSMG_coarse_solve', 11)
551 if (allocated(this%amg_solver)) then
552 call this%amg_solver%solve(this%grids(1)%e%x, this%wf%x, &
553 this%grids(1)%dof%size())
554 else
555 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, &
556 this%wf%x, &
557 this%grids(1)%dof%size(), &
558 this%grids(1)%coef, &
559 this%grids(1)%bc_projector, &
560 this%grids(1)%gs_h, this%niter)
561 end if
562 call profiler_end_region('HSMG_coarse_solve', 11)
563 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x,&
564 this%grids(1)%dof%size())
565 call profiler_end_region('HSMG_coarse_grid', 10)
566 end if
567 !$omp end parallel
568
569 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
570 this%msh%nelv, this%grids(2)%Xh)
571 call device_add2(this%grids(2)%e%x_d, this%w_d, this%grids(2)%dof%size())
572
573 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
574 this%msh%nelv, this%grids(3)%Xh)
575 call device_add2(z_d, this%w_d, this%grids(3)%dof%size())
576 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), &
577 gs_op_add, this%gs_event)
578 call device_event_sync(this%gs_event)
579 call device_col2(z_d, this%grids(3)%coef%mult_d, &
580 this%grids(3)%dof%size())
581 else
582 !We should not work with the input
583 call copy(this%r, r, n)
584
585 !OVERLAPPING Schwarz exchange and solve
586 call this%grids(3)%schwarz%compute(z, this%r)
587 ! DOWNWARD Leg of V-cycle, we are pretty hardcoded here but w/e
588 call col2(this%r, this%grids(3)%coef%mult, &
589 this%grids(3)%dof%size())
590 !Restrict to middle level
591 call this%interp_fine_mid%map(this%w, this%r, &
592 this%msh%nelv, this%grids(2)%Xh)
593 call this%grids(2)%gs_h%op(this%w, this%grids(2)%dof%size(), gs_op_add)
594 !OVERLAPPING Schwarz exchange and solve
595 call this%grids(2)%schwarz%compute(this%grids(2)%e%x, this%w)
596 call col2(this%w, this%grids(2)%coef%mult, this%grids(2)%dof%size())
597 !restrict residual to crs
598 call this%interp_mid_crs%map(this%r, this%w, &
599 this%msh%nelv, this%grids(1)%Xh)
600 !Crs solve
601
602 call this%grids(1)%gs_h%op(this%r, this%grids(1)%dof%size(), gs_op_add)
603 call this%grids(1)%bc_projector%apply(this%r, this%grids(1)%dof%size())
604
605 call profiler_start_region('HSMG_coarse-solve', 11)
606 if (allocated(this%amg_solver)) then
607 call this%amg_solver%solve(this%grids(1)%e%x, this%r, &
608 this%grids(1)%dof%size())
609 else
610 crs_info = this%crs_solver%solve(this%Ax, this%grids(1)%e, this%r, &
611 this%grids(1)%dof%size(), &
612 this%grids(1)%coef, &
613 this%grids(1)%bc_projector, &
614 this%grids(1)%gs_h, this%niter)
615 end if
616 call profiler_end_region('HSMG_coarse-solve', 11)
617
618 call this%grids(1)%bc_projector%apply(this%grids(1)%e%x, &
619 this%grids(1)%dof%size())
620
621
622 call this%interp_mid_crs%map(this%w, this%grids(1)%e%x, &
623 this%msh%nelv, this%grids(2)%Xh)
624 call add2(this%grids(2)%e%x, this%w, this%grids(2)%dof%size())
625
626 call this%interp_fine_mid%map(this%w, this%grids(2)%e%x, &
627 this%msh%nelv, this%grids(3)%Xh)
628 call add2(z, this%w, this%grids(3)%dof%size())
629 call this%grids(3)%gs_h%op(z, this%grids(3)%dof%size(), gs_op_add)
630 call col2(z, this%grids(3)%coef%mult, this%grids(3)%dof%size())
631
632 end if
633 call profiler_end_region('HSMG_solve', 8)
634 end subroutine hsmg_solve
635end 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:132
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, json)
Definition pc_hsmg.f90:150
subroutine hsmg_solve(this, z, r, n)
The h1mg preconditioner from Nek5000.
Definition pc_hsmg.f90:489
subroutine hsmg_set_h(this)
Definition pc_hsmg.f90:374
subroutine hsmg_free(this)
Definition pc_hsmg.f90:411
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:389
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:92
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:44
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:136
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:47
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:44
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