Neko 1.99.7
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
40 use field, only : field_t
41 use coefs, only : coef_t
42 use mesh, only : mesh_t
43 use space, only : space_t
44 use gather_scatter, only : gs_t, gs_op_add
45 use bc_list, only : bc_list_t
46 use schwarz, only : schwarz_t
53 use, intrinsic :: iso_c_binding, only : c_ptr, c_int, &
54 c_null_ptr, c_associated
55 implicit none
56 private
57
59 type, public, extends(ksp_t) :: cheby_device_t
60 real(kind=rp), allocatable :: d(:)
61 real(kind=rp), allocatable :: w(:)
62 real(kind=rp), allocatable :: r(:)
63 type(c_ptr) :: d_d = c_null_ptr
64 type(c_ptr) :: w_d = c_null_ptr
65 type(c_ptr) :: r_d = c_null_ptr
66 type(c_ptr) :: gs_event = c_null_ptr
67 real(kind=rp) :: tha, dlt
68 integer :: power_its = 150
70 integer :: power_its_refresh = 20
72 logical :: warm_start_eigs = .false.
74 logical :: eigs_computed = .false.
76 real(kind=rp), allocatable :: ev(:)
77 type(c_ptr) :: ev_d = c_null_ptr
78 logical :: recompute_eigs = .true.
79 logical :: zero_initial_guess = .false.
80 type(schwarz_t), pointer :: schwarz => null()
81 contains
82 procedure, pass(this) :: init => cheby_device_init
83 procedure, pass(this) :: free => cheby_device_free
84 procedure, pass(this) :: solve => cheby_device_impl
85 procedure, pass(this) :: solve_coupled => cheby_device_solve_coupled
86 end type cheby_device_t
87
88#ifdef HAVE_HIP
89 interface
90 subroutine hip_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
91 bind(c, name = 'hip_cheby_part1')
92 use, intrinsic :: iso_c_binding
93 import c_rp
94 implicit none
95 type(c_ptr), value :: d_d, x_d, strm
96 real(c_rp) :: inv_tha
97 integer(c_int) :: n
98 end subroutine hip_cheby_device_part1
99 end interface
100
101 interface
102 subroutine hip_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
103 bind(c, name = 'hip_cheby_part2')
104 use, intrinsic :: iso_c_binding
105 import c_rp
106 implicit none
107 type(c_ptr), value :: d_d, w_d, x_d, strm
108 real(c_rp) :: tmp1, tmp2
109 integer(c_int) :: n
110 end subroutine hip_cheby_device_part2
111 end interface
112#elif HAVE_CUDA
113 interface
114 subroutine cuda_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
115 bind(c, name = 'cuda_cheby_part1')
116 use, intrinsic :: iso_c_binding
117 import c_rp
118 implicit none
119 type(c_ptr), value :: d_d, x_d, strm
120 real(c_rp) :: inv_tha
121 integer(c_int) :: n
122 end subroutine cuda_cheby_device_part1
123 end interface
124
125 interface
126 subroutine cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
127 bind(c, name = 'cuda_cheby_part2')
128 use, intrinsic :: iso_c_binding
129 import c_rp
130 implicit none
131 type(c_ptr), value :: d_d, w_d, x_d, strm
132 real(c_rp) :: tmp1, tmp2
133 integer(c_int) :: n
134 end subroutine cuda_cheby_device_part2
135 end interface
136#elif HAVE_METAL
137 interface
138 subroutine metal_cheby_device_part1(d_d, x_d, inv_tha, n, strm) &
139 bind(c, name = 'metal_cheby_part1')
140 use, intrinsic :: iso_c_binding
141 import c_rp
142 implicit none
143 type(c_ptr), value :: d_d, x_d, strm
144 real(c_rp) :: inv_tha
145 integer(c_int) :: n
146 end subroutine metal_cheby_device_part1
147 end interface
148
149 interface
150 subroutine metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, strm) &
151 bind(c, name = 'metal_cheby_part2')
152 use, intrinsic :: iso_c_binding
153 import c_rp
154 implicit none
155 type(c_ptr), value :: d_d, w_d, x_d, strm
156 real(c_rp) :: tmp1, tmp2
157 integer(c_int) :: n
158 end subroutine metal_cheby_device_part2
159 end interface
160#endif
161
162contains
163 subroutine cheby_device_part1(d_d, x_d, inv_tha, n)
164 type(c_ptr) :: d_d, x_d
165 real(c_rp) :: inv_tha
166 integer(c_int) :: n
167#ifdef HAVE_HIP
168 call hip_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
169#elif HAVE_CUDA
170 call cuda_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
171#elif HAVE_METAL
172 call metal_cheby_device_part1(d_d, x_d, inv_tha, n, glb_cmd_queue)
173#else !Fallback to device_math for missing device kernels
174 call device_cmult( d_d, inv_tha, n)
175 call device_add2( x_d, d_d, n)
176#endif
177 end subroutine cheby_device_part1
178
179
180
181 subroutine cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n)
182 type(c_ptr) :: d_d, w_d, x_d
183 real(c_rp) :: tmp1, tmp2
184 integer(c_int) :: n
185#ifdef HAVE_HIP
186 call hip_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
187#elif HAVE_CUDA
188 call cuda_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
189#elif HAVE_METAL
190 call metal_cheby_device_part2(d_d, w_d, x_d, tmp1, tmp2, n, glb_cmd_queue)
191#else !Fallback to device_math for missing device kernels
192 call device_cmult( d_d, tmp1, n)
193 call device_add2s2( d_d, w_d, tmp2, n)
194 call device_add2( x_d, d_d, n)
195#endif
196 end subroutine cheby_device_part2
197
199 subroutine cheby_device_init(this, n, max_iter, M, rel_tol, abs_tol, monitor)
200 class(cheby_device_t), intent(inout), target :: this
201 integer, intent(in) :: max_iter
202 class(pc_t), optional, intent(in), target :: M
203 integer, intent(in) :: n
204 real(kind=rp), optional, intent(in) :: rel_tol
205 real(kind=rp), optional, intent(in) :: abs_tol
206 logical, optional, intent(in) :: monitor
207
208 call this%free()
209 allocate(this%d(n))
210 allocate(this%w(n))
211 allocate(this%r(n))
212
213 call device_map(this%d, this%d_d, n)
214 call device_map(this%w, this%w_d, n)
215 call device_map(this%r, this%r_d, n)
216
217 if (present(m)) then
218 this%M => m
219 end if
220
221 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
222 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
223 else if (present(rel_tol) .and. present(abs_tol)) then
224 call this%ksp_init(max_iter, rel_tol, abs_tol)
225 else if (present(monitor) .and. present(abs_tol)) then
226 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
227 else if (present(rel_tol) .and. present(monitor)) then
228 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
229 else if (present(rel_tol)) then
230 call this%ksp_init(max_iter, rel_tol = rel_tol)
231 else if (present(abs_tol)) then
232 call this%ksp_init(max_iter, abs_tol = abs_tol)
233 else if (present(monitor)) then
234 call this%ksp_init(max_iter, monitor = monitor)
235 else
236 call this%ksp_init(max_iter)
237 end if
238
239 call device_event_create(this%gs_event, 2)
240
241 end subroutine cheby_device_init
242
243 subroutine cheby_device_free(this)
244 class(cheby_device_t), intent(inout) :: this
245
246 call this%ksp_free()
247
248 if (allocated(this%d)) then
249 if (c_associated(this%d_d)) then
250 call device_unmap(this%d, this%d_d)
251 end if
252 deallocate(this%d)
253 end if
254
255 if (allocated(this%w)) then
256 if (c_associated(this%w_d)) then
257 call device_unmap(this%w, this%w_d)
258 end if
259 deallocate(this%w)
260 end if
261
262 if (allocated(this%r)) then
263 if (c_associated(this%r_d)) then
264 call device_unmap(this%r, this%r_d)
265 end if
266 deallocate(this%r)
267 end if
268
269 if (allocated(this%ev)) then
270 if (c_associated(this%ev_d)) then
271 call device_unmap(this%ev, this%ev_d)
272 end if
273 deallocate(this%ev)
274 end if
275
276 nullify(this%M)
277
278 if (c_associated(this%gs_event)) then
279 call device_event_destroy(this%gs_event)
280 end if
281
282 end subroutine cheby_device_free
283
284 subroutine cheby_device_power(this, Ax, x, n, coef, blst, gs_h)
285 class(cheby_device_t), intent(inout) :: this
286 class(ax_t), intent(in) :: Ax
287 type(field_t), intent(inout) :: x
288 integer, intent(in) :: n
289 type(coef_t), intent(inout) :: coef
290 type(bc_list_t), intent(inout) :: blst
291 type(gs_t), intent(inout) :: gs_h
292 real(kind=rp) :: lam, b, a, rn
293 real(kind=rp) :: boost = 1.1_rp
294 real(kind=rp) :: lam_factor = 30.0_rp
295 real(kind=rp) :: wtw, dtw, dtd
296 integer, allocatable :: fixed_seed(:), saved_seed(:)
297 integer :: i, rnd_n, its
298 logical :: warm
299
300 call profiler_start_region('cheby_power')
301 warm = this%warm_start_eigs .and. this%eigs_computed .and. &
302 c_associated(this%ev_d)
303 associate(w => this%w, w_d => this%w_d, d => this%d, d_d => this%d_d)
304
305 if (warm) then
306 its = this%power_its_refresh
307 call device_copy(d_d, this%ev_d, n)
308 else
309 its = this%power_its
310
311 ! Save current random seed and set a fixed seed
312 call random_seed(size = rnd_n)
313 allocate(saved_seed(rnd_n))
314 allocate(fixed_seed(rnd_n))
315 fixed_seed = 3901
316 call random_seed(get = saved_seed)
317 call random_seed(put = fixed_seed)
318
319 do i = 1, n
320 call random_number(rn)
321 d(i) = rn + 10.0_rp
322 end do
323 call device_memcpy(d, d_d, n, host_to_device, sync = .true.)
324
325 ! Restore saved random seed
326 call random_seed(put = saved_seed)
327
328 call gs_h%op(d, n, gs_op_add, this%gs_event)
329 call blst%apply(d, n)
330 end if
331
332 !Power method to get lambda max
333 do i = 1, its
334 call ax%compute(w, d, coef, x%msh, x%Xh)
335 call gs_h%op(w, n, gs_op_add, this%gs_event)
336 call blst%apply(w, n)
337 if (associated(this%schwarz)) then
338 call this%schwarz%compute(this%r, w)
339 call device_copy(w_d, this%r_d, n)
340 else
341 call this%M%solve(this%r, w, n)
342 call device_copy(w_d, this%r_d, n)
343 end if
344
345 wtw = device_glsc3(w_d, coef%mult_d, w_d, n)
346 call device_cmult2(d_d, w_d, 1.0_rp/sqrt(wtw), n)
347 call blst%apply(d, n)
348 end do
349
350 call ax%compute(w, d, coef, x%msh, x%Xh)
351 call gs_h%op(w, n, gs_op_add, this%gs_event)
352 call blst%apply(w, n)
353 if (associated(this%schwarz)) then
354 call this%schwarz%compute(this%r, w)
355 call device_copy(w_d, this%r_d, n)
356 else
357 call this%M%solve(this%r, w, n)
358 call device_copy(w_d, this%r_d, n)
359 end if
360
361 dtw = device_glsc3(d_d, coef%mult_d, w_d, n)
362 dtd = device_glsc3(d_d, coef%mult_d, d_d, n)
363 lam = dtw / dtd
364 b = lam * boost
365 a = lam / lam_factor
366 this%tha = (b+a)/2.0_rp
367 this%dlt = (b-a)/2.0_rp
368
369 if (this%warm_start_eigs) then
370 if (.not. c_associated(this%ev_d)) then
371 allocate(this%ev(size(this%d)))
372 call device_map(this%ev, this%ev_d, size(this%d))
373 end if
374 call device_copy(this%ev_d, d_d, n)
375 end if
376 this%eigs_computed = .true.
377
378 this%recompute_eigs = .false.
379 end associate
380 call profiler_end_region('cheby_power')
381 end subroutine cheby_device_power
382
384 function cheby_device_solve(this, Ax, x, f, n, coef, blst, gs_h, niter) &
385 result(ksp_results)
386 class(cheby_device_t), intent(inout) :: this
387 class(ax_t), intent(in) :: ax
388 type(field_t), intent(inout) :: x
389 integer, intent(in) :: n
390 real(kind=rp), dimension(n), intent(in) :: f
391 type(coef_t), intent(inout) :: coef
392 type(bc_list_t), intent(inout) :: blst
393 type(gs_t), intent(inout) :: gs_h
394 type(ksp_monitor_t) :: ksp_results
395 integer, optional, intent(in) :: niter
396 integer :: iter, max_iter
397 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
398 type(c_ptr) :: f_d
399
400 f_d = device_get_ptr(f)
401
402 if (this%recompute_eigs) then
403 call cheby_device_power(this, ax, x, n, coef, blst, gs_h)
404 end if
405
406 if (present(niter)) then
407 max_iter = niter
408 else
409 max_iter = this%max_iter
410 end if
411 norm_fac = 1.0_rp / sqrt(coef%volume)
412
413 associate( w => this%w, r => this%r, d => this%d, &
414 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
415 ! calculate residual
416 call device_copy(r_d, f_d, n)
417 call ax%compute(w, x%x, coef, x%msh, x%Xh)
418 call gs_h%op(w, n, gs_op_add, this%gs_event)
419 call blst%apply(w, n)
420 call device_sub2(r_d, w_d, n)
421
422 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
423 rnorm = sqrt(rtr) * norm_fac
424 ksp_results%res_start = rnorm
425 ksp_results%res_final = rnorm
426 ksp_results%iter = 0
427
428 ! First iteration
429 call this%M%solve(w, r, n)
430 call device_copy(d_d, w_d, n)
431 a = 2.0_rp / this%tha
432 call device_add2s2(x%x_d, d_d, a, n)! x = x + a*d
433
434 ! Rest of the iterations
435 do iter = 2, max_iter
436 ! calculate residual
437 call device_copy(r_d, f_d, n)
438 call ax%compute(w, x%x, coef, x%msh, x%Xh)
439 call gs_h%op(w, n, gs_op_add, this%gs_event)
440 call blst%apply(w, n)
441 call device_sub2(r_d, w_d, n)
442
443 call this%M%solve(w, r, n)
444
445 if (iter .eq. 2) then
446 b = 0.5_rp * (this%dlt * a)**2
447 else
448 b = (this%dlt * a / 2.0_rp)**2
449 end if
450 a = 1.0_rp/(this%tha - b/a)
451 call device_add2s1(d_d, w_d, b, n)! d = w + b*d
452
453 call device_add2s2(x%x_d, d_d, a, n)! x = x + a*d
454 end do
455
456 ! calculate residual
457 call device_copy(r_d, f_d, n)
458 call ax%compute(w, x%x, coef, x%msh, x%Xh)
459 call gs_h%op(w, n, gs_op_add, this%gs_event)
460 call blst%apply(w, n)
461 call device_sub2(r_d, w_d, n)
462 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
463 rnorm = sqrt(rtr) * norm_fac
464
465
466 ksp_results%res_final = rnorm
467 ksp_results%iter = iter
468 ksp_results%converged = this%is_converged(iter, rnorm)
469 end associate
470 end function cheby_device_solve
471
473 function cheby_device_impl(this, Ax, x, f, n, coef, blst, gs_h, niter) &
474 result(ksp_results)
475 class(cheby_device_t), intent(inout) :: this
476 class(ax_t), intent(in) :: ax
477 type(field_t), intent(inout) :: x
478 integer, intent(in) :: n
479 real(kind=rp), dimension(n), intent(in) :: f
480 type(coef_t), intent(inout) :: coef
481 type(bc_list_t), intent(inout) :: blst
482 type(gs_t), intent(inout) :: gs_h
483 type(ksp_monitor_t) :: ksp_results
484 integer, optional, intent(in) :: niter
485 integer :: iter, max_iter
486 real(kind=rp) :: a, b, rtr, rnorm, norm_fac
487 real(kind=rp) :: rhok, rhokp1, sig1, tmp1, tmp2
488 type(c_ptr) :: f_d
489
490 f_d = device_get_ptr(f)
491
492 if (this%recompute_eigs) then
493 call cheby_device_power(this, ax, x, n, coef, blst, gs_h)
494 end if
495
496 if (present(niter)) then
497 max_iter = niter
498 else
499 max_iter = this%max_iter
500 end if
501 norm_fac = 1.0_rp / sqrt(coef%volume)
502
503 associate( w => this%w, r => this%r, d => this%d, &
504 w_d => this%w_d, r_d => this%r_d, d_d => this%d_d)
505 ! calculate residual
506 if (.not.this%zero_initial_guess) then
507 call ax%compute(w, x%x, coef, x%msh, x%Xh)
508 call gs_h%op(w, n, gs_op_add, this%gs_event)
509 call blst%apply(w, n)
510 call device_sub3(r_d, f_d, w_d, n)
511 else
512 call device_copy(r_d, f_d, n)
513 this%zero_initial_guess = .false.
514 end if
515
516 ! First iteration
517 if (associated(this%schwarz)) then
518 call this%schwarz%compute(d, r)
519 else
520 call this%M%solve(d, r, n)
521 end if
522
523 tmp1 = 1.0_rp / this%tha
524 call cheby_device_part1(d_d, x%x_d, tmp1, n)
525
526 sig1 = this%tha / this%dlt
527 rhok = 1.0_rp / sig1
528
529 ! Rest of the iterations
530 do iter = 2, max_iter
531 rhokp1 = 1.0_rp / (2.0_rp * sig1 - rhok)
532 tmp1 = rhokp1 * rhok
533 tmp2 = 2.0_rp * rhokp1 / this%dlt
534 rhok = rhokp1
535 ! calculate residual
536 call ax%compute(w, x%x, coef, x%msh, x%Xh)
537 call gs_h%op(w, n, gs_op_add, this%gs_event)
538 call blst%apply(w, n)
539 call device_sub3(r_d, f_d, w_d, n)
540
541 if (associated(this%schwarz)) then
542 call this%schwarz%compute(w, r)
543 else
544 call this%M%solve(w, r, n)
545 end if
546
547 call cheby_device_part2(d_d, w_d, x%x_d, tmp1, tmp2, n)
548
549 end do
550
551 end associate
552 end function cheby_device_impl
553
555 function cheby_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
556 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
557 class(cheby_device_t), intent(inout) :: this
558 class(ax_t), intent(in) :: ax
559 type(field_t), intent(inout) :: x
560 type(field_t), intent(inout) :: y
561 type(field_t), intent(inout) :: z
562 integer, intent(in) :: n
563 real(kind=rp), dimension(n), intent(in) :: fx
564 real(kind=rp), dimension(n), intent(in) :: fy
565 real(kind=rp), dimension(n), intent(in) :: fz
566 type(coef_t), intent(inout) :: coef
567 type(bc_list_t), intent(inout) :: blstx
568 type(bc_list_t), intent(inout) :: blsty
569 type(bc_list_t), intent(inout) :: blstz
570 type(gs_t), intent(inout) :: gs_h
571 type(ksp_monitor_t), dimension(3) :: ksp_results
572 integer, optional, intent(in) :: niter
573
574 ksp_results(1) = this%solve(ax, x, fx, n, coef, blstx, gs_h, niter)
575 ksp_results(2) = this%solve(ax, y, fy, n, coef, blsty, gs_h, niter)
576 ksp_results(3) = this%solve(ax, z, fz, n, coef, blstz, gs_h, niter)
577
578 end function cheby_device_solve_coupled
579
580end 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
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
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