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