Neko 1.99.6
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
cheby_device.F90
Go to the documentation of this file.
1! Copyright (c) 2024-2025, The Neko Authors
2! All rights reserved.
3!
4! Redistribution and use in source and binary forms, with or without
5! modification, are permitted provided that the following conditions
6! are met:
7!
8! * Redistributions of source code must retain the above copyright
9! notice, this list of conditions and the following disclaimer.
10!
11! * Redistributions in binary form must reproduce the above
12! copyright notice, this list of conditions and the following
13! disclaimer in the documentation and/or other materials provided
14! with the distribution.
15!
16! * Neither the name of the authors nor the names of its
17! contributors may be used to endorse or promote products derived
18! from this software without specific prior written permission.
19!
20! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
21! "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
22! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
23! FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
24! COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
25! INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
26! BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
27! LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
28! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
29! LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
30! ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
31! POSSIBILITY OF SUCH DAMAGE.
32!
35 use krylov, only : ksp_t, ksp_monitor_t
36 use precon, only : pc_t
37 use ax_product, only : ax_t
38 use num_types, only : rp, c_rp
39 use field, only : field_t
40 use coefs, only : coef_t
41 use mesh, only : mesh_t
42 use space, only : space_t
43 use gather_scatter, only : gs_t, gs_op_add
44 use bc_list, only : bc_list_t
45 use schwarz, only : schwarz_t
52 use, intrinsic :: iso_c_binding, only : c_ptr, c_int, &
53 c_null_ptr, c_associated
54 implicit none
55 private
56
58 type, public, extends(ksp_t) :: cheby_device_t
59 real(kind=rp), allocatable :: d(:)
60 real(kind=rp), allocatable :: w(:)
61 real(kind=rp), allocatable :: r(:)
62 type(c_ptr) :: d_d = c_null_ptr
63 type(c_ptr) :: w_d = c_null_ptr
64 type(c_ptr) :: r_d = c_null_ptr
65 type(c_ptr) :: gs_event = c_null_ptr
66 real(kind=rp) :: tha, dlt
67 integer :: power_its = 150
68 logical :: recompute_eigs = .true.
69 logical :: zero_initial_guess = .false.
70 type(schwarz_t), pointer :: schwarz => null()
71 contains
72 procedure, pass(this) :: init => cheby_device_init
73 procedure, pass(this) :: free => cheby_device_free
74 procedure, pass(this) :: solve => cheby_device_impl
75 procedure, pass(this) :: solve_coupled => cheby_device_solve_coupled
76 end type cheby_device_t
77
78#ifdef HAVE_HIP
79 interface
80 subroutine hip_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
81 bind(c, name = 'hip_cheby_part1')
82 use, intrinsic :: iso_c_binding
83 import c_rp
84 implicit none
85 type(c_ptr), value :: d_d, x_d, strm
86 real(c_rp) :: inv_tha
87 integer(c_int) :: n
88 end subroutine hip_cheby_device_part1
89 end interface
90
91 interface
92 subroutine hip_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
93 bind(c, name = 'hip_cheby_part2')
94 use, intrinsic :: iso_c_binding
95 import c_rp
96 implicit none
97 type(c_ptr), value :: d_d, w_d, x_d, strm
98 real(c_rp) :: tmp1, tmp2
99 integer(c_int) :: n
100 end subroutine hip_cheby_device_part2
101 end interface
102#elif HAVE_CUDA
103 interface
104 subroutine cuda_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
105 bind(c, name = 'cuda_cheby_part1')
106 use, intrinsic :: iso_c_binding
107 import c_rp
108 implicit none
109 type(c_ptr), value :: d_d, x_d, strm
110 real(c_rp) :: inv_tha
111 integer(c_int) :: n
112 end subroutine cuda_cheby_device_part1
113 end interface
114
115 interface
116 subroutine cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
117 bind(c, name = 'cuda_cheby_part2')
118 use, intrinsic :: iso_c_binding
119 import c_rp
120 implicit none
121 type(c_ptr), value :: d_d, w_d, x_d, strm
122 real(c_rp) :: tmp1, tmp2
123 integer(c_int) :: n
124 end subroutine cuda_cheby_device_part2
125 end interface
126#elif HAVE_METAL
127 interface
128 subroutine metal_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
129 bind(c, name = 'metal_cheby_part1')
130 use, intrinsic :: iso_c_binding
131 import c_rp
132 implicit none
133 type(c_ptr), value :: d_d, x_d, strm
134 real(c_rp) :: inv_tha
135 integer(c_int) :: n
136 end subroutine metal_cheby_device_part1
137 end interface
138
139 interface
140 subroutine metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
141 bind(c, name = 'metal_cheby_part2')
142 use, intrinsic :: iso_c_binding
143 import c_rp
144 implicit none
145 type(c_ptr), value :: d_d, w_d, x_d, strm
146 real(c_rp) :: tmp1, tmp2
147 integer(c_int) :: n
148 end subroutine metal_cheby_device_part2
149 end interface
150#endif
151
152contains
153 subroutine cheby_device_part1(d_d, x_d, inv_tha, n)
154 type(c_ptr) :: d_d, x_d
155 real(c_rp) :: inv_tha
156 integer(c_int) :: n
157#ifdef HAVE_HIP
158 call hip_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
159#elif HAVE_CUDA
160 call cuda_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
161#elif HAVE_METAL
162 call metal_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
163#else !Fallback to device_math for missing device kernels
164 call device_cmult( d_d, inv_tha, n)
165 call device_add2( x_d, d_d, n)
166#endif
167 end subroutine cheby_device_part1
168
169
170
171 subroutine cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n)
172 type(c_ptr) :: d_d, w_d, x_d
173 real(c_rp) :: tmp1, tmp2
174 integer(c_int) :: n
175#ifdef HAVE_HIP
176 call hip_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
177#elif HAVE_CUDA
178 call cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
179#elif HAVE_METAL
180 call metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
181#else !Fallback to device_math for missing device kernels
182 call device_cmult( d_d, tmp1, n)
183 call device_add2s2( d_d, w_d, tmp2, n)
184 call device_add2( x_d, d_d, n)
185#endif
186 end subroutine cheby_device_part2
187
189 subroutine cheby_device_init(this, n, max_iter, M, rel_tol, abs_tol, monitor)
190 class(cheby_device_t), intent(inout), target :: this
191 integer, intent(in) :: max_iter
192 class(pc_t), optional, intent(in), target :: M
193 integer, intent(in) :: n
194 real(kind=rp), optional, intent(in) :: rel_tol
195 real(kind=rp), optional, intent(in) :: abs_tol
196 logical, optional, intent(in) :: monitor
197
198 call this%free()
199 allocate(this%d(n))
200 allocate(this%w(n))
201 allocate(this%r(n))
202
203 call device_map(this%d, this%d_d, n)
204 call device_map(this%w, this%w_d, n)
205 call device_map(this%r, this%r_d, n)
206
207 if (present(m)) then
208 this%M => m
209 end if
210
211 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
212 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
213 else if (present(rel_tol) .and. present(abs_tol)) then
214 call this%ksp_init(max_iter, rel_tol, abs_tol)
215 else if (present(monitor) .and. present(abs_tol)) then
216 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
217 else if (present(rel_tol) .and. present(monitor)) then
218 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
219 else if (present(rel_tol)) then
220 call this%ksp_init(max_iter, rel_tol = rel_tol)
221 else if (present(abs_tol)) then
222 call this%ksp_init(max_iter, abs_tol = abs_tol)
223 else if (present(monitor)) then
224 call this%ksp_init(max_iter, monitor = monitor)
225 else
226 call this%ksp_init(max_iter)
227 end if
228
229 call device_event_create(this%gs_event, 2)
230
231 end subroutine cheby_device_init
232
233 subroutine cheby_device_free(this)
234 class(cheby_device_t), intent(inout) :: this
235
236 call this%ksp_free()
237
238 if (allocated(this%d)) then
239 if (c_associated(this%d_d)) then
240 call device_unmap(this%d, this%d_d)
241 end if
242 deallocate(this%d)
243 end if
244
245 if (allocated(this%w)) then
246 if (c_associated(this%w_d)) then
247 call device_unmap(this%w, this%w_d)
248 end if
249 deallocate(this%w)
250 end if
251
252 if (allocated(this%r)) then
253 if (c_associated(this%r_d)) then
254 call device_unmap(this%r, this%r_d)
255 end if
256 deallocate(this%r)
257 end if
258
259 nullify(this%M)
260
261 if (c_associated(this%gs_event)) then
262 call device_event_destroy(this%gs_event)
263 end if
264
265 end subroutine cheby_device_free
266
267 subroutine cheby_device_power(this, Ax, x, n, coef, blst, gs_h)
268 class(cheby_device_t), intent(inout) :: this
269 class(ax_t), intent(in) :: Ax
270 type(field_t), intent(inout) :: x
271 integer, intent(in) :: n
272 type(coef_t), intent(inout) :: coef
273 type(bc_list_t), intent(inout) :: blst
274 type(gs_t), intent(inout) :: gs_h
275 real(kind=rp) :: lam, b, a, rn
276 real(kind=rp) :: boost = 1.1_rp
277 real(kind=rp) :: lam_factor = 30.0_rp
278 real(kind=rp) :: wtw, dtw, dtd
279 integer, allocatable :: fixed_seed(:), saved_seed(:)
280 integer :: i, rnd_n
281
282 associate(w => this%w, w_d => this%w_d, d => this%d, d_d => this%d_d)
283
284 ! Save current random seed and set a fixed seed
285 call random_seed( size = rnd_n )
286 allocate(saved_seed(rnd_n))
287 allocate(fixed_seed(rnd_n))
288 fixed_seed = 3901
289 call random_seed( get = saved_seed )
290 call random_seed( put = fixed_seed )
291
292 do i = 1, n
293 call random_number(rn)
294 d(i) = rn + 10.0_rp
295 end do
296 call device_memcpy(d, d_d, n, host_to_device, sync = .true.)
297
298 ! Restore saved random seed
299 call random_seed( put = saved_seed )
300
301 call gs_h%op(d, n, gs_op_add, this%gs_event)
302 call blst%apply(d, n)
303
304 !Power method to get lamba max
305 do i = 1, this%power_its
306 call ax%compute(w, d, coef, x%msh, x%Xh)
307 call gs_h%op(w, n, gs_op_add, this%gs_event)
308 call blst%apply(w, n)
309 if (associated(this%schwarz)) then
310 call this%schwarz%compute(this%r, w)
311 call device_copy(w_d, this%r_d, n)
312 else
313 call this%M%solve(this%r, w, n)
314 call device_copy(w_d, this%r_d, n)
315 end if
316
317 wtw = device_glsc3(w_d, coef%mult_d, w_d, n)
318 call device_cmult2(d_d, w_d, 1.0_rp/sqrt(wtw), n)
319 call blst%apply(d, n)
320 end do
321
322 call ax%compute(w, d, coef, x%msh, x%Xh)
323 call gs_h%op(w, n, gs_op_add, this%gs_event)
324 call blst%apply(w, n)
325 if (associated(this%schwarz)) then
326 call this%schwarz%compute(this%r, w)
327 call device_copy(w_d, this%r_d, n)
328 else
329 call this%M%solve(this%r, w, n)
330 call device_copy(w_d, this%r_d, n)
331 end if
332
333 dtw = device_glsc3(d_d, coef%mult_d, w_d, n)
334 dtd = device_glsc3(d_d, coef%mult_d, d_d, n)
335 lam = dtw / dtd
336 b = lam * boost
337 a = lam / lam_factor
338 this%tha = (b+a)/2.0_rp
339 this%dlt = (b-a)/2.0_rp
340
341 this%recompute_eigs = .false.
342 end associate
343 end subroutine cheby_device_power
344
346 function cheby_device_solve(this, Ax, x, f, n, coef, blst, gs_h, niter) &
347 result(ksp_results)
348 class(cheby_device_t), intent(inout) :: this
349 class(ax_t), intent(in) :: ax
350 type(field_t), intent(inout) :: x
351 integer, intent(in) :: n
352 real(kind=rp), dimension(n), intent(in) :: f
353 type(coef_t), intent(inout) :: coef
354 type(bc_list_t), intent(inout) :: blst
355 type(gs_t), intent(inout) :: gs_h
356 type(ksp_monitor_t) :: ksp_results
357 integer, optional, intent(in) :: niter
358 integer :: iter, max_iter
359 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
360 type(c_ptr) :: f_d
361
362 f_d = device_get_ptr(f)
363
364 if (this%recompute_eigs) then
365 call cheby_device_power(this, ax, x, n, coef, blst, gs_h)
366 end if
367
368 if (present(niter)) then
369 max_iter = niter
370 else
371 max_iter = this%max_iter
372 end if
373 norm_fac = 1.0_rp / sqrt(coef%volume)
374
375 associate( w => this%w, r => this%r, d => this%d, &
376 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
377 ! calculate residual
378 call device_copy(r_d, f_d, n)
379 call ax%compute(w, x%x, coef, x%msh, x%Xh)
380 call gs_h%op(w, n, gs_op_add, this%gs_event)
381 call blst%apply(w, n)
382 call device_sub2(r_d, w_d, n)
383
384 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
385 rnorm = sqrt(rtr) * norm_fac
386 ksp_results%res_start = rnorm
387 ksp_results%res_final = rnorm
388 ksp_results%iter = 0
389
390 ! First iteration
391 call this%M%solve(w, r, n)
392 call device_copy(d_d, w_d, n)
393 a = 2.0_rp / this%tha
394 call device_add2s2(x%x_d, d_d, a, n)! x = x + a*d
395
396 ! Rest of the iterations
397 do iter = 2, max_iter
398 ! calculate residual
399 call device_copy(r_d, f_d, n)
400 call ax%compute(w, x%x, coef, x%msh, x%Xh)
401 call gs_h%op(w, n, gs_op_add, this%gs_event)
402 call blst%apply(w, n)
403 call device_sub2(r_d, w_d, n)
404
405 call this%M%solve(w, r, n)
406
407 if (iter .eq. 2) then
408 b = 0.5_rp * (this%dlt * a)**2
409 else
410 b = (this%dlt * a / 2.0_rp)**2
411 end if
412 a = 1.0_rp/(this%tha - b/a)
413 call device_add2s1(d_d, w_d, b, n)! d = w + b*d
414
415 call device_add2s2(x%x_d, d_d, a, n)! x = x + a*d
416 end do
417
418 ! calculate residual
419 call device_copy(r_d, f_d, n)
420 call ax%compute(w, x%x, coef, x%msh, x%Xh)
421 call gs_h%op(w, n, gs_op_add, this%gs_event)
422 call blst%apply(w, n)
423 call device_sub2(r_d, w_d, n)
424 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
425 rnorm = sqrt(rtr) * norm_fac
426
427
428 ksp_results%res_final = rnorm
429 ksp_results%iter = iter
430 ksp_results%converged = this%is_converged(iter, rnorm)
431 end associate
432 end function cheby_device_solve
433
435 function cheby_device_impl(this, Ax, x, f, n, coef, blst, gs_h, niter) &
436 result(ksp_results)
437 class(cheby_device_t), intent(inout) :: this
438 class(ax_t), intent(in) :: ax
439 type(field_t), intent(inout) :: x
440 integer, intent(in) :: n
441 real(kind=rp), dimension(n), intent(in) :: f
442 type(coef_t), intent(inout) :: coef
443 type(bc_list_t), intent(inout) :: blst
444 type(gs_t), intent(inout) :: gs_h
445 type(ksp_monitor_t) :: ksp_results
446 integer, optional, intent(in) :: niter
447 integer :: iter, max_iter
448 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
449 real(kind=rp) :: rhok, rhokp1, sig1, tmp1, tmp2
450 type(c_ptr) :: f_d
451
452 f_d = device_get_ptr(f)
453
454 if (this%recompute_eigs) then
455 call cheby_device_power(this, ax, x, n, coef, blst, gs_h)
456 end if
457
458 if (present(niter)) then
459 max_iter = niter
460 else
461 max_iter = this%max_iter
462 end if
463 norm_fac = 1.0_rp / sqrt(coef%volume)
464
465 associate( w => this%w, r => this%r, d => this%d, &
466 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
467 ! calculate residual
468 if (.not.this%zero_initial_guess) then
469 call ax%compute(w, x%x, coef, x%msh, x%Xh)
470 call gs_h%op(w, n, gs_op_add, this%gs_event)
471 call blst%apply(w, n)
472 call device_sub3(r_d, f_d, w_d, n)
473 else
474 call device_copy(r_d, f_d, n)
475 this%zero_initial_guess = .false.
476 end if
477
478 ! First iteration
479 if (associated(this%schwarz)) then
480 call this%schwarz%compute(d, r)
481 else
482 call this%M%solve(d, r, n)
483 end if
484
485 tmp1 = 1.0_rp / this%tha
486 call cheby_device_part1(d_d, x%x_d, tmp1, n)
487
488 sig1 = this%tha / this%dlt
489 rhok = 1.0_rp / sig1
490
491 ! Rest of the iterations
492 do iter = 2, max_iter
493 rhokp1 = 1.0_rp / (2.0_rp * sig1 - rhok)
494 tmp1 = rhokp1 * rhok
495 tmp2 = 2.0_rp * rhokp1 / this%dlt
496 rhok = rhokp1
497 ! calculate residual
498 call ax%compute(w, x%x, coef, x%msh, x%Xh)
499 call gs_h%op(w, n, gs_op_add, this%gs_event)
500 call blst%apply(w, n)
501 call device_sub3(r_d, f_d, w_d, n)
502
503 if (associated(this%schwarz)) then
504 call this%schwarz%compute(w, r)
505 else
506 call this%M%solve(w, r, n)
507 end if
508
509 call cheby_device_part2(d_d, w_d, x%x_d, tmp1, tmp2, n)
510
511 end do
512
513 end associate
514 end function cheby_device_impl
515
517 function cheby_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
518 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
519 class(cheby_device_t), intent(inout) :: this
520 class(ax_t), intent(in) :: ax
521 type(field_t), intent(inout) :: x
522 type(field_t), intent(inout) :: y
523 type(field_t), intent(inout) :: z
524 integer, intent(in) :: n
525 real(kind=rp), dimension(n), intent(in) :: fx
526 real(kind=rp), dimension(n), intent(in) :: fy
527 real(kind=rp), dimension(n), intent(in) :: fz
528 type(coef_t), intent(inout) :: coef
529 type(bc_list_t), intent(inout) :: blstx
530 type(bc_list_t), intent(inout) :: blsty
531 type(bc_list_t), intent(inout) :: blstz
532 type(gs_t), intent(inout) :: gs_h
533 type(ksp_monitor_t), dimension(3) :: ksp_results
534 integer, optional, intent(in) :: niter
535
536 ksp_results(1) = this%solve(ax, x, fx, n, coef, blstx, gs_h, niter)
537 ksp_results(2) = this%solve(ax, y, fy, n, coef, blsty, gs_h, niter)
538 ksp_results(3) = this%solve(ax, z, fz, n, coef, blstz, gs_h, niter)
539
540 end function cheby_device_solve_coupled
541
542end module cheby_device
__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
Copy data between host and device (or device and device)
Definition device.F90:72
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Defines a Matrix-vector product.
Definition ax.f90:34
Defines a list of bc_t.
Definition bc_list.f90:34
Chebyshev preconditioner.
subroutine cheby_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a standard solver.
subroutine cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n)
type(ksp_monitor_t) function, dimension(3) cheby_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, blstx, blsty, blstz, gs_h, niter)
Standard Cheby_Deviceshev coupled solve.
type(ksp_monitor_t) function cheby_device_solve(this, ax, x, f, n, coef, blst, gs_h, niter)
A chebyshev preconditioner.
subroutine cheby_device_free(this)
subroutine cheby_device_part1(d_d, x_d, inv_tha, n)
subroutine cheby_device_power(this, ax, x, n, coef, blst, gs_h)
type(ksp_monitor_t) function cheby_device_impl(this, ax, x, f, n, coef, blst, gs_h, niter)
A chebyshev preconditioner.
Coefficients.
Definition coef.f90:34
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_sub3(a_d, b_d, c_d, n, strm)
Vector subtraction .
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_cmult(a_d, c, n, strm)
Multiplication by constant c .
subroutine, public device_sub2(a_d, b_d, n, strm)
Vector substraction .
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
real(kind=rp) function, public device_glsc3(a_d, b_d, c_d, n, strm)
Weighted inner product .
subroutine, public device_cmult2(a_d, b_d, c, n, strm)
Multiplication by constant c .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
integer, parameter, public host_to_device
Definition device.F90:48
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1623
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
Definition device.F90:52
subroutine, public device_event_create(event, flags)
Create a device event queue.
Definition device.F90:1589
Defines a field.
Definition field.f90:34
Gather-scatter.
Implements the base abstract type for Krylov solvers plus helper types.
Definition krylov.f90:34
Defines a mesh.
Definition mesh.f90:34
integer, parameter, public c_rp
Definition num_types.f90:13
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Krylov preconditioner.
Definition precon.f90:34
Overlapping schwarz solves.
Definition schwarz.f90:61
Defines a function space.
Definition space.f90:34
Base type for a matrix-vector product providing .
Definition ax.f90:43
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Defines a Chebyshev preconditioner.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:63
Gather-scatter kernel.
Type for storing initial and final residuals in a Krylov solver.
Definition krylov.f90:56
Base abstract type for a canonical Krylov method, solving .
Definition krylov.f90:73
Defines a canonical Krylov preconditioner.
Definition precon.f90:40
The function space for the SEM solution fields.
Definition space.f90:64