Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
bicgstab_device.F90
Go to the documentation of this file.
1! Copyright (c) 2026, 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 num_types, only : rp, xp, c_rp, c_xp
36 use krylov, only : ksp_t, ksp_monitor_t
37 use precon, only : pc_t
38 use ax_product, only : ax_t
39 use field, only : field_t
40 use coefs, only : coef_t
41 use gather_scatter, only : gs_t, gs_op_add
45 use math, only : neko_eps
51 use utils, only : neko_error
53 use mpi_f08, only : mpi_allreduce, mpi_in_place, mpi_sum
54 use, intrinsic :: ieee_arithmetic, only : ieee_is_finite
55 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated, &
56 c_int
57 implicit none
58 private
59
71 type, public, extends(ksp_t) :: bicgstab_device_t
73 real(kind=rp), allocatable :: p(:)
75 real(kind=rp), allocatable :: p_hat(:)
77 real(kind=rp), allocatable :: r(:)
79 real(kind=rp), allocatable :: s(:)
81 real(kind=rp), allocatable :: s_hat(:)
83 real(kind=rp), allocatable :: t(:)
85 real(kind=rp), allocatable :: v(:)
86 type(c_ptr) :: p_d = c_null_ptr
87 type(c_ptr) :: p_hat_d = c_null_ptr
88 type(c_ptr) :: r_d = c_null_ptr
89 type(c_ptr) :: s_d = c_null_ptr
90 type(c_ptr) :: s_hat_d = c_null_ptr
91 type(c_ptr) :: t_d = c_null_ptr
92 type(c_ptr) :: v_d = c_null_ptr
93 type(c_ptr) :: gs_event = c_null_ptr
94 contains
96 procedure, pass(this) :: init => bicgstab_device_init
98 procedure, pass(this) :: free => bicgstab_device_free
100 procedure, pass(this) :: solve => bicgstab_device_solve
102 procedure, pass(this) :: solve_coupled => bicgstab_device_solve_coupled
103 end type bicgstab_device_t
104
105#if HAVE_HIP
106 interface
107 subroutine hip_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n) &
108 bind(c, name = 'hip_bicgstab_update_p')
109 use, intrinsic :: iso_c_binding
110 import c_rp
111 implicit none
112 type(c_ptr), value :: p_d, r_d, v_d
113 real(c_rp) :: beta, omega
114 integer(c_int) :: n
115 end subroutine hip_bicgstab_update_p
116 end interface
117
118 interface
119 subroutine hip_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n) &
120 bind(c, name = 'hip_bicgstab_product_and_norm')
121 use, intrinsic :: iso_c_binding
122 import c_xp
123 implicit none
124 type(c_ptr), value :: a_d, b_d, mult_d
125 real(c_xp) :: res(2)
126 integer(c_int) :: n
127 end subroutine hip_bicgstab_product_and_norm
128 end interface
129
130 interface
131 real(c_xp) function hip_bicgstab_part1(s_d, r_d, v_d, mult_d, &
132 alpha, n) bind(c, name = 'hip_bicgstab_part1')
133 use, intrinsic :: iso_c_binding
134 import c_rp, c_xp
135 implicit none
136 type(c_ptr), value :: s_d, r_d, v_d, mult_d
137 real(c_rp) :: alpha
138 integer(c_int) :: n
139 end function hip_bicgstab_part1
140 end interface
141
142 interface
143 subroutine hip_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, &
144 t_d, f_d, mult_d, alpha, omega, res, n) &
145 bind(c, name = 'hip_bicgstab_part2')
146 use, intrinsic :: iso_c_binding
147 import c_rp, c_xp
148 implicit none
149 type(c_ptr), value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
150 type(c_ptr), value :: mult_d
151 real(c_rp) :: alpha, omega
152 real(c_xp) :: res(2)
153 integer(c_int) :: n
154 end subroutine hip_bicgstab_part2
155 end interface
156#elif HAVE_CUDA
157 interface
158 subroutine cuda_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n) &
159 bind(c, name = 'cuda_bicgstab_update_p')
160 use, intrinsic :: iso_c_binding
161 import c_rp
162 implicit none
163 type(c_ptr), value :: p_d, r_d, v_d
164 real(c_rp) :: beta, omega
165 integer(c_int) :: n
166 end subroutine cuda_bicgstab_update_p
167 end interface
168
169 interface
170 subroutine cuda_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n) &
171 bind(c, name = 'cuda_bicgstab_product_and_norm')
172 use, intrinsic :: iso_c_binding
173 import c_xp
174 implicit none
175 type(c_ptr), value :: a_d, b_d, mult_d
176 real(c_xp) :: res(2)
177 integer(c_int) :: n
178 end subroutine cuda_bicgstab_product_and_norm
179 end interface
180
181 interface
182 real(c_xp) function cuda_bicgstab_part1(s_d, r_d, v_d, mult_d, &
183 alpha, n) bind(c, name = 'cuda_bicgstab_part1')
184 use, intrinsic :: iso_c_binding
185 import c_rp, c_xp
186 implicit none
187 type(c_ptr), value :: s_d, r_d, v_d, mult_d
188 real(c_rp) :: alpha
189 integer(c_int) :: n
190 end function cuda_bicgstab_part1
191 end interface
192
193 interface
194 subroutine cuda_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, &
195 t_d, f_d, mult_d, alpha, omega, res, n) &
196 bind(c, name = 'cuda_bicgstab_part2')
197 use, intrinsic :: iso_c_binding
198 import c_rp, c_xp
199 implicit none
200 type(c_ptr), value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
201 type(c_ptr), value :: mult_d
202 real(c_rp) :: alpha, omega
203 real(c_xp) :: res(2)
204 integer(c_int) :: n
205 end subroutine cuda_bicgstab_part2
206 end interface
207#elif HAVE_OPENCL
208 interface
209 subroutine opencl_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
210 strm) bind(c, name = 'opencl_bicgstab_update_p')
211 use, intrinsic :: iso_c_binding
212 import c_rp
213 implicit none
214 type(c_ptr), value :: p_d, r_d, v_d, strm
215 real(c_rp) :: beta, omega
216 integer(c_int) :: n
217 end subroutine opencl_bicgstab_update_p
218 end interface
219
220 interface
221 subroutine opencl_bicgstab_product_and_norm(a_d, b_d, mult_d, res, &
222 n, strm) bind(c, name = 'opencl_bicgstab_product_and_norm')
223 use, intrinsic :: iso_c_binding
224 import c_xp
225 implicit none
226 type(c_ptr), value :: a_d, b_d, mult_d, strm
227 real(c_xp) :: res(2)
228 integer(c_int) :: n
230 end interface
231
232 interface
233 real(c_xp) function opencl_bicgstab_part1(s_d, r_d, v_d, mult_d, &
234 alpha, n, strm) bind(c, name = 'opencl_bicgstab_part1')
235 use, intrinsic :: iso_c_binding
236 import c_rp, c_xp
237 implicit none
238 type(c_ptr), value :: s_d, r_d, v_d, mult_d, strm
239 real(c_rp) :: alpha
240 integer(c_int) :: n
241 end function opencl_bicgstab_part1
242 end interface
243
244 interface
245 subroutine opencl_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, &
246 t_d, f_d, mult_d, alpha, omega, res, n, strm) &
247 bind(c, name = 'opencl_bicgstab_part2')
248 use, intrinsic :: iso_c_binding
249 import c_rp, c_xp
250 implicit none
251 type(c_ptr), value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
252 type(c_ptr), value :: mult_d, strm
253 real(c_rp) :: alpha, omega
254 real(c_xp) :: res(2)
255 integer(c_int) :: n
256 end subroutine opencl_bicgstab_part2
257 end interface
258#elif HAVE_METAL
259 interface
260 subroutine metal_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
261 strm) bind(c, name = 'metal_bicgstab_update_p')
262 use, intrinsic :: iso_c_binding
263 import c_rp
264 implicit none
265 type(c_ptr), value :: p_d, r_d, v_d, strm
266 real(c_rp) :: beta, omega
267 integer(c_int) :: n
268 end subroutine metal_bicgstab_update_p
269 end interface
270
271 interface
272 subroutine metal_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n, &
273 strm) bind(c, name = 'metal_bicgstab_product_and_norm')
274 use, intrinsic :: iso_c_binding
275 import c_rp
276 implicit none
277 type(c_ptr), value :: a_d, b_d, mult_d, strm
278 real(c_rp) :: res(2)
279 integer(c_int) :: n
280 end subroutine metal_bicgstab_product_and_norm
281 end interface
282
283 interface
284 real(c_rp) function metal_bicgstab_part1(s_d, r_d, v_d, mult_d, &
285 alpha, n, strm) bind(c, name = 'metal_bicgstab_part1')
286 use, intrinsic :: iso_c_binding
287 import c_rp
288 implicit none
289 type(c_ptr), value :: s_d, r_d, v_d, mult_d, strm
290 real(c_rp) :: alpha
291 integer(c_int) :: n
292 end function metal_bicgstab_part1
293 end interface
294
295 interface
296 subroutine metal_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, &
297 t_d, f_d, mult_d, alpha, omega, res, n, strm) &
298 bind(c, name = 'metal_bicgstab_part2')
299 use, intrinsic :: iso_c_binding
300 import c_rp
301 implicit none
302 type(c_ptr), value :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d
303 type(c_ptr), value :: mult_d, strm
304 real(c_rp) :: alpha, omega, res(2)
305 integer(c_int) :: n
306 end subroutine metal_bicgstab_part2
307 end interface
308#endif
309
310contains
311
319 subroutine device_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
320 type(c_ptr) :: p_d, r_d, v_d
321 real(kind=rp) :: beta, omega
322 integer :: n
323
324#if HAVE_HIP
325 call hip_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
326#elif HAVE_CUDA
327 call cuda_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
328#elif HAVE_OPENCL
329 call opencl_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
330 glb_cmd_queue)
331#elif HAVE_METAL
332 call metal_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n, &
333 glb_cmd_queue)
334#else
335 call neko_error('No device backend configured')
336#endif
337
338 end subroutine device_bicgstab_update_p
339
351 subroutine device_bicgstab_product_and_norm(product, norm_squared, &
352 a_d, b_d, mult_d, n)
353 real(kind=rp), intent(out) :: product
354 real(kind=rp), intent(out) :: norm_squared
355 type(c_ptr) :: a_d, b_d, mult_d
356 integer :: n
357 real(kind=xp) :: res_xp(2)
358 integer :: ierr
359#if HAVE_METAL
360 real(kind=rp) :: res(2)
361#endif
362
363 res_xp = 0.0_xp
364#if HAVE_HIP
365 call hip_bicgstab_product_and_norm(a_d, b_d, mult_d, res_xp, n)
366#elif HAVE_CUDA
367 call cuda_bicgstab_product_and_norm(a_d, b_d, mult_d, res_xp, n)
368#elif HAVE_OPENCL
369 call opencl_bicgstab_product_and_norm(a_d, b_d, mult_d, res_xp, n, &
370 glb_cmd_queue)
371#elif HAVE_METAL
372 call metal_bicgstab_product_and_norm(a_d, b_d, mult_d, res, n, &
373 glb_cmd_queue)
374 res_xp = real(res, kind=xp)
375#else
376 call neko_error('No device backend configured')
377#endif
378
379#ifndef HAVE_DEVICE_MPI
380 if (pe_size .gt. 1) then
381 call mpi_allreduce(mpi_in_place, res_xp, 2, &
382 mpi_extra_precision, mpi_sum, neko_comm, ierr)
383 end if
384#endif
385 product = real(res_xp(1), kind=rp)
386 norm_squared = real(res_xp(2), kind=rp)
387
389
398 function device_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n) &
399 result(res)
400 type(c_ptr) :: s_d, r_d, v_d, mult_d
401 real(kind=rp) :: alpha
402 integer :: n
403 real(kind=rp) :: res
404 real(kind=xp) :: res_xp
405 integer :: ierr
406
407 res_xp = 0.0_xp
408#if HAVE_HIP
409 res_xp = hip_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n)
410#elif HAVE_CUDA
411 res_xp = cuda_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n)
412#elif HAVE_OPENCL
413 res_xp = opencl_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n, &
414 glb_cmd_queue)
415#elif HAVE_METAL
416 res = metal_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n, &
417 glb_cmd_queue)
418 res_xp = real(res, kind=xp)
419#else
420 call neko_error('No device backend configured')
421#endif
422
423#ifndef HAVE_DEVICE_MPI
424 if (pe_size .gt. 1) then
425 call mpi_allreduce(mpi_in_place, res_xp, 1, &
426 mpi_extra_precision, mpi_sum, neko_comm, ierr)
427 end if
428#endif
429 res = real(res_xp, kind=rp)
430
431 end function device_bicgstab_part1
432
452 subroutine device_bicgstab_part2(rtr, rho, x_d, r_d, p_hat_d, s_hat_d, &
453 s_d, t_d, f_d, mult_d, alpha, omega, n)
454 real(kind=rp), intent(out) :: rtr
455 real(kind=rp), intent(out) :: rho
456 type(c_ptr) :: x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, mult_d
457 real(kind=rp) :: alpha, omega
458 integer :: n
459 real(kind=xp) :: res_xp(2)
460 integer :: ierr
461#if HAVE_METAL
462 real(kind=rp) :: res(2)
463#endif
464
465 res_xp = 0.0_xp
466#if HAVE_HIP
467 call hip_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, &
468 mult_d, alpha, omega, res_xp, n)
469#elif HAVE_CUDA
470 call cuda_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, &
471 mult_d, alpha, omega, res_xp, n)
472#elif HAVE_OPENCL
473 call opencl_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, &
474 mult_d, alpha, omega, res_xp, n, glb_cmd_queue)
475#elif HAVE_METAL
476 call metal_bicgstab_part2(x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, &
477 mult_d, alpha, omega, res, n, glb_cmd_queue)
478 res_xp = real(res, kind=xp)
479#else
480 call neko_error('No device backend configured')
481#endif
482
483#ifndef HAVE_DEVICE_MPI
484 if (pe_size .gt. 1) then
485 call mpi_allreduce(mpi_in_place, res_xp, 2, &
486 mpi_extra_precision, mpi_sum, neko_comm, ierr)
487 end if
488#endif
489 rtr = real(res_xp(1), kind=rp)
490 rho = real(res_xp(2), kind=rp)
491
492 end subroutine device_bicgstab_part2
493
502 subroutine bicgstab_device_init(this, n, max_iter, M, rel_tol, abs_tol, &
503 monitor)
504 class(bicgstab_device_t), target, intent(inout) :: this
505 class(pc_t), optional, intent(in), target :: M
506 integer, intent(in) :: n
507 integer, intent(in) :: max_iter
508 real(kind=rp), optional, intent(in) :: rel_tol
509 real(kind=rp), optional, intent(in) :: abs_tol
510 logical, optional, intent(in) :: monitor
511
512 call this%free()
513
514 allocate(this%p(n))
515 allocate(this%p_hat(n))
516 allocate(this%r(n))
517 allocate(this%s(n))
518 allocate(this%s_hat(n))
519 allocate(this%t(n))
520 allocate(this%v(n))
521
522 call device_map(this%p, this%p_d, n)
523 call device_map(this%p_hat, this%p_hat_d, n)
524 call device_map(this%r, this%r_d, n)
525 call device_map(this%s, this%s_d, n)
526 call device_map(this%s_hat, this%s_hat_d, n)
527 call device_map(this%t, this%t_d, n)
528 call device_map(this%v, this%v_d, n)
529
530 if (present(m)) then
531 this%M => m
532 end if
533
534 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
535 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
536 else if (present(rel_tol) .and. present(abs_tol)) then
537 call this%ksp_init(max_iter, rel_tol, abs_tol)
538 else if (present(monitor) .and. present(abs_tol)) then
539 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
540 else if (present(rel_tol) .and. present(monitor)) then
541 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
542 else if (present(rel_tol)) then
543 call this%ksp_init(max_iter, rel_tol = rel_tol)
544 else if (present(abs_tol)) then
545 call this%ksp_init(max_iter, abs_tol = abs_tol)
546 else if (present(monitor)) then
547 call this%ksp_init(max_iter, monitor = monitor)
548 else
549 call this%ksp_init(max_iter)
550 end if
551
552 call device_event_create(this%gs_event, 2)
553
554 end subroutine bicgstab_device_init
555
557 subroutine bicgstab_device_free(this)
558 class(bicgstab_device_t), intent(inout) :: this
559
560 call this%ksp_free()
561
562 if (allocated(this%v)) then
563 if (c_associated(this%v_d)) then
564 call device_unmap(this%v, this%v_d)
565 end if
566 deallocate(this%v)
567 end if
568
569 if (allocated(this%r)) then
570 if (c_associated(this%r_d)) then
571 call device_unmap(this%r, this%r_d)
572 end if
573 deallocate(this%r)
574 end if
575
576 if (allocated(this%t)) then
577 if (c_associated(this%t_d)) then
578 call device_unmap(this%t, this%t_d)
579 end if
580 deallocate(this%t)
581 end if
582
583 if (allocated(this%p)) then
584 if (c_associated(this%p_d)) then
585 call device_unmap(this%p, this%p_d)
586 end if
587 deallocate(this%p)
588 end if
589
590 if (allocated(this%p_hat)) then
591 if (c_associated(this%p_hat_d)) then
592 call device_unmap(this%p_hat, this%p_hat_d)
593 end if
594 deallocate(this%p_hat)
595 end if
596
597 if (allocated(this%s)) then
598 if (c_associated(this%s_d)) then
599 call device_unmap(this%s, this%s_d)
600 end if
601 deallocate(this%s)
602 end if
603
604 if (allocated(this%s_hat)) then
605 if (c_associated(this%s_hat_d)) then
606 call device_unmap(this%s_hat, this%s_hat_d)
607 end if
608 deallocate(this%s_hat)
609 end if
610
611 nullify(this%M)
612
613 if (c_associated(this%gs_event)) then
614 call device_event_destroy(this%gs_event)
615 end if
616
617 end subroutine bicgstab_device_free
618
632 function bicgstab_device_solve(this, Ax, x, f, n, coef, bc_projector, &
633 gs_h, niter) result(ksp_results)
634 class(bicgstab_device_t), intent(inout) :: this
635 class(ax_t), intent(in) :: ax
636 type(field_t), intent(inout) :: x
637 integer, intent(in) :: n
638 real(kind=rp), dimension(n), intent(in) :: f
639 type(coef_t), intent(inout) :: coef
640 class(scalar_bc_projector_t), intent(inout) :: bc_projector
641 type(gs_t), intent(inout) :: gs_h
642 type(ksp_monitor_t) :: ksp_results
643 integer, optional, intent(in) :: niter
644 integer :: iter, max_iter
645 real(kind=rp) :: rnorm, rtr, norm_fac, gamma
646 real(kind=rp) :: r_norm, s_norm, shadow_norm, t_norm, v_norm
647 ! s^T s, f^T v, v^T v, s^T t, t^T t
648 real(kind=rp) :: sts, ftv, vtv, stt, ttt
649 real(kind=rp) :: beta, alpha, omega, rho_1, rho_2, rho_next
650 type(c_ptr) :: f_d
651
652 f_d = device_get_ptr(f)
653
654 if (present(niter)) then
655 max_iter = niter
656 else
657 max_iter = this%max_iter
658 end if
659 norm_fac = 1.0_rp / sqrt(coef%volume)
660
661 associate(r_d => this%r_d, t_d => this%t_d, s_d => this%s_d, &
662 v_d => this%v_d, p_d => this%p_d, s_hat_d => this%s_hat_d, &
663 p_hat_d => this%p_hat_d, mult_d => coef%mult_d)
664
665 call device_rzero(x%x_d, n)
666 call device_copy(r_d, f_d, n)
667 rtr = device_glsc3(r_d, mult_d, r_d, n)
668
669 ! The implementation deliberately starts from x = 0, so f is both the
670 ! initial residual and the fixed shadow residual used by BiCGStab.
671 r_norm = bicgstab_device_sqrt(rtr, 'initial residual')
672 shadow_norm = r_norm
673 rnorm = r_norm * norm_fac
674 gamma = rnorm * this%rel_tol
675 ksp_results%res_start = rnorm
676 ksp_results%res_final = rnorm
677 ksp_results%iter = 0
678
679 ! Avoid entering the recurrence if the zero initial guess already meets
680 ! either stopping criterion. Apart from saving work, this prevents a
681 ! small right-hand side from being mistaken for a rho breakdown.
682 if (r_norm .le. 0.0_rp .or. rnorm .lt. this%abs_tol .or. &
683 rnorm .lt. gamma) then
684 ksp_results%converged = .true.
685 return
686 end if
687
688 ! Every later rho comes out of part 2 together with the residual norm.
689 ! On the first iteration r is still f, so rho is the norm just reduced.
690 rho_1 = rtr
691
692 call this%monitor_start('BiCGStab')
693 do iter = 1, max_iter
694
695 ! BiCGStab breaks down if the residual becomes orthogonal to the
696 ! shadow residual. Test the inner product relative to the norms of
697 ! both vectors so that uniformly scaling the system does not change
698 ! the decision.
699 call bicgstab_device_check_inner_product(rho_1, shadow_norm, &
700 r_norm, 'rho inner product')
701
702 if (iter .eq. 1) then
703 call device_copy(p_d, r_d, n)
704 else
705 beta = (rho_1 / rho_2) * (alpha / omega)
706 if (.not. ieee_is_finite(beta)) then
707 call neko_error('BiCGStab failure: non-finite beta')
708 end if
709 call device_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
710 end if
711
712 call this%M%solve(this%p_hat, this%p, n)
713 call ax%compute(this%v, this%p_hat, coef, x%msh, x%Xh)
714 call gs_h%op(this%v, n, gs_op_add, this%gs_event)
715 call device_event_sync(this%gs_event)
716 call bc_projector%apply(this%v, n)
717
718 ! The alpha denominator is another BiCG breakdown point. Reducing it
719 ! together with ||v|| permits a scale-aware orthogonality check
720 ! without an additional global synchronisation.
721 call device_bicgstab_product_and_norm(ftv, vtv, f_d, v_d, &
722 mult_d, n)
723 v_norm = bicgstab_device_sqrt(vtv, 'operator result v')
724 call bicgstab_device_check_inner_product(ftv, shadow_norm, &
725 v_norm, 'alpha denominator')
726 alpha = rho_1 / ftv
727 if (.not. ieee_is_finite(alpha)) then
728 call neko_error('BiCGStab failure: non-finite alpha')
729 end if
730
731 sts = device_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n)
732
733 s_norm = bicgstab_device_sqrt(sts, 'intermediate residual')
734 rnorm = s_norm * norm_fac
735 if (rnorm .lt. this%abs_tol .or. rnorm .lt. gamma) then
736 call device_add2s2(x%x_d, p_hat_d, alpha, n)
737 call this%monitor_iter(iter, rnorm)
738 exit
739 end if
740
741 call this%M%solve(this%s_hat, this%s, n)
742 call ax%compute(this%t, this%s_hat, coef, x%msh, x%Xh)
743 call gs_h%op(this%t, n, gs_op_add, this%gs_event)
744 call device_event_sync(this%gs_event)
745 call bc_projector%apply(this%t, n)
746
747 call device_bicgstab_product_and_norm(stt, ttt, s_d, t_d, &
748 mult_d, n)
749 t_norm = bicgstab_device_sqrt(ttt, 'operator result t')
750 if (t_norm .le. 0.0_rp) then
751 call neko_error('BiCGStab breakdown: zero omega denominator')
752 end if
753 if (.not. ieee_is_finite(stt)) then
754 call neko_error('BiCGStab failure: non-finite omega numerator')
755 end if
756 omega = stt / ttt
757 if (.not. ieee_is_finite(omega)) then
758 call neko_error('BiCGStab failure: non-finite omega')
759 end if
760
761 call device_bicgstab_part2(rtr, rho_next, x%x_d, r_d, p_hat_d, &
762 s_hat_d, s_d, t_d, f_d, mult_d, alpha, omega, n)
763
764 r_norm = bicgstab_device_sqrt(rtr, 'recursive residual')
765 rnorm = r_norm * norm_fac
766 call this%monitor_iter(iter, rnorm)
767 if (rnorm .lt. this%abs_tol .or. rnorm .lt. gamma) then
768 exit
769 end if
770
771 ! A negative omega is valid. Breakdown occurs only when the local
772 ! minimal-residual step stagnates, i.e. when t and s are numerically
773 ! orthogonal.
774 call bicgstab_device_check_inner_product(stt, t_norm, s_norm, &
775 'omega numerator')
776 rho_2 = rho_1
777 rho_1 = rho_next
778
779 end do
780 end associate
781 call this%monitor_stop()
782 ksp_results%res_final = rnorm
783 ksp_results%iter = iter
784 ksp_results%converged = this%is_converged(iter, rnorm)
785
786 end function bicgstab_device_solve
787
796 subroutine bicgstab_device_check_inner_product(inner_product, norm_a, &
797 norm_b, quantity)
798 real(kind=rp), intent(in) :: inner_product
799 real(kind=rp), intent(in) :: norm_a
800 real(kind=rp), intent(in) :: norm_b
801 character(len=*), intent(in) :: quantity
802 real(kind=rp) :: large_norm, small_norm
803
804 if (.not. ieee_is_finite(inner_product)) then
805 call neko_error('BiCGStab failure: non-finite ' // trim(quantity))
806 end if
807
808 large_norm = max(norm_a, norm_b)
809 small_norm = min(norm_a, norm_b)
810 if (large_norm .le. 0.0_rp .or. &
811 abs(inner_product) / large_norm .le. neko_eps * small_norm) then
812 call neko_error('BiCGStab breakdown: near-zero ' // trim(quantity))
813 end if
814
816
821 function bicgstab_device_sqrt(value, quantity) result(root)
822 real(kind=rp), intent(in) :: value
823 character(len=*), intent(in) :: quantity
824 real(kind=rp) :: root
825
826 if (.not. ieee_is_finite(value) .or. value .lt. 0.0_rp) then
827 call neko_error('BiCGStab failure: invalid ' // trim(quantity) // &
828 ' norm')
829 end if
830 root = sqrt(value)
831
832 end function bicgstab_device_sqrt
833
851 function bicgstab_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
852 n, coef, bc_projector, gs_h, niter) result(ksp_results)
853 class(bicgstab_device_t), intent(inout) :: this
854 class(ax_t), intent(in) :: ax
855 type(field_t), intent(inout) :: x
856 type(field_t), intent(inout) :: y
857 type(field_t), intent(inout) :: z
858 integer, intent(in) :: n
859 real(kind=rp), dimension(n), intent(in) :: fx
860 real(kind=rp), dimension(n), intent(in) :: fy
861 real(kind=rp), dimension(n), intent(in) :: fz
862 type(coef_t), intent(inout) :: coef
863 class(vector_bc_projector_t), intent(inout) :: bc_projector
864 type(gs_t), intent(inout) :: gs_h
865 type(ksp_monitor_t), dimension(3) :: ksp_results
866 integer, optional, intent(in) :: niter
867 type(scalar_bc_projector_t), pointer :: bc_x, bc_y, bc_z
868
869 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
870 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
871 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
872 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
873
875
876end module bicgstab_device
real_xp opencl_bicgstab_part1(void *s, void *r, void *v, void *mult, real *alpha, int *n, cl_command_queue cmd_queue)
Definition bicgstab.c:160
void opencl_bicgstab_part2(void *x, void *r, void *p_hat, void *s_hat, void *s, void *t, void *f, void *mult, real *alpha, real *omega, real_xp *res, int *n, cl_command_queue cmd_queue)
Definition bicgstab.c:218
void opencl_bicgstab_product_and_norm(void *a, void *b, void *mult, real_xp *res, int *n, cl_command_queue cmd_queue)
Definition bicgstab.c:104
void opencl_bicgstab_update_p(void *p, void *r, void *v, real *beta, real *omega, int *n, cl_command_queue cmd_queue)
Definition bicgstab.c:67
void cuda_bicgstab_product_and_norm(void *a, void *b, void *mult, real_xp *res, int *n)
void cuda_bicgstab_update_p(void *p, void *r, void *v, real *beta, real *omega, int *n)
void cuda_bicgstab_part2(void *x, void *r, void *p_hat, void *s_hat, void *s, void *t, void *f, void *mult, real *alpha, real *omega, real_xp *res, int *n)
real_xp cuda_bicgstab_part1(void *s, void *r, void *v, void *mult, real *alpha, int *n)
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
double real
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
Defines a Matrix-vector product.
Definition ax.f90:34
Provides a device implementation of the BiCGStab method.
subroutine device_bicgstab_update_p(p_d, r_d, v_d, beta, omega, n)
Search direction update .
subroutine bicgstab_device_check_inner_product(inner_product, norm_a, norm_b, quantity)
Check an inner product for a BiCGStab breakdown.
type(ksp_monitor_t) function bicgstab_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Solve a linear system with the device BiCGStab method.
subroutine device_bicgstab_product_and_norm(product, norm_squared, a_d, b_d, mult_d, n)
Weighted inner product and squared norm in one reduction.
type(ksp_monitor_t) function, dimension(3) bicgstab_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Solve three independent systems with the device BiCGStab method.
real(kind=rp) function device_bicgstab_part1(s_d, r_d, v_d, mult_d, alpha, n)
BiCGStab part 1, .
subroutine bicgstab_device_free(this)
Free a device BiCGStab solver.
subroutine bicgstab_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device BiCGStab solver.
real(kind=rp) function bicgstab_device_sqrt(value, quantity)
Return the square root of a valid squared norm.
subroutine device_bicgstab_part2(rtr, rho, x_d, r_d, p_hat_d, s_hat_d, s_d, t_d, f_d, mult_d, alpha, omega, n)
BiCGStab part 2, and .
Coefficients.
Definition coef.f90:34
Definition comm.F90:1
integer, public pe_size
MPI size of communicator.
Definition comm.F90:62
type(mpi_comm), public neko_comm
MPI communicator.
Definition comm.F90:46
type(mpi_datatype), public mpi_extra_precision
Definition comm.F90:55
subroutine, public device_add2s2(a_d, b_d, c1, n, strm)
Vector addition with scalar multiplication (multiplication on first argument)
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
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 .
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
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
Definition math.f90:60
real(kind=rp), parameter, public neko_eps
Machine epsilon .
Definition math.f90:70
integer, parameter, public c_xp
Definition num_types.f90:17
integer, parameter, public xp
Definition num_types.f90:16
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
Implements scalar_projector_t.
Utilities.
Definition utils.f90:35
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
Device implementation of the right-preconditioned BiCGStab method.
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.
Abstract type for resolving vector boundary conditions.
#define max(a, b)
Definition tensor.cu:40