Neko 1.99.6
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
pipecg_device.F90
Go to the documentation of this file.
1! Copyright (c) 2021-2024, 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!
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 gather_scatter, only : gs_t, gs_op_add
42 use bc_list, only : bc_list_t
43 use math, only : glsc3, rzero, copy, abscmp
49 use utils, only : neko_error
51 use mpi_f08, only : mpi_iallreduce, mpi_status, &
52 mpi_sum, mpi_in_place, mpi_request, mpi_wait
53 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, &
54 c_associated, c_size_t, c_sizeof, c_int, c_loc
55 implicit none
56 private
57
58 integer, parameter :: device_pipecg_p_space = 10
59
61 type, public, extends(ksp_t) :: pipecg_device_t
62 real(kind=rp), allocatable :: p(:)
63 real(kind=rp), allocatable :: q(:)
64 real(kind=rp), allocatable :: r(:)
65 real(kind=rp), allocatable :: s(:)
66 real(kind=rp), allocatable :: u(:,:)
67 real(kind=rp), allocatable :: w(:)
68 real(kind=rp), allocatable :: z(:)
69 real(kind=rp), allocatable :: mi(:)
70 real(kind=rp), allocatable :: ni(:)
71 real(kind=rp), allocatable :: alpha(:)
72 real(kind=rp), allocatable :: beta(:)
73 type(c_ptr) :: p_d = c_null_ptr
74 type(c_ptr) :: q_d = c_null_ptr
75 type(c_ptr) :: r_d = c_null_ptr
76 type(c_ptr) :: s_d = c_null_ptr
77 type(c_ptr) :: u_d_d = c_null_ptr
78 type(c_ptr) :: w_d = c_null_ptr
79 type(c_ptr) :: z_d = c_null_ptr
80 type(c_ptr) :: mi_d = c_null_ptr
81 type(c_ptr) :: ni_d = c_null_ptr
82 type(c_ptr) :: alpha_d = c_null_ptr
83 type(c_ptr) :: beta_d = c_null_ptr
84 type(c_ptr), allocatable :: u_d(:)
85 type(c_ptr) :: gs_event = c_null_ptr
86 contains
87 procedure, pass(this) :: init => pipecg_device_init
88 procedure, pass(this) :: free => pipecg_device_free
89 procedure, pass(this) :: solve => pipecg_device_solve
90 procedure, pass(this) :: solve_coupled => pipecg_device_solve_coupled
91 end type pipecg_device_t
92
93#ifdef HAVE_CUDA
94 interface
95 subroutine cuda_pipecg_vecops(p_d, q_d, r_d, s_d, u_d1, u_d2, &
96 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n) &
97 bind(c, name = 'cuda_pipecg_vecops')
98 use, intrinsic :: iso_c_binding
99 import c_rp
100 implicit none
101 type(c_ptr), value :: p_d, q_d, r_d, s_d, u_d1, u_d2
102 type(c_ptr), value :: w_d, ni_d, mi_d, z_d, mult_d
103 integer(c_int) :: n
104 real(c_rp) :: alpha, beta, reduction(3)
105 end subroutine cuda_pipecg_vecops
106 end interface
107
108 interface
109 subroutine cuda_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, &
110 p_cur, p_space, n) &
111 bind(c, name = 'cuda_cg_update_xp')
112 use, intrinsic :: iso_c_binding
113 implicit none
114 type(c_ptr), value :: x_d, p_d, u_d_d, alpha, beta
115 integer(c_int) :: p_cur, n, p_space
116 end subroutine cuda_cg_update_xp
117 end interface
118#elif HAVE_HIP
119 interface
120 subroutine hip_pipecg_vecops(p_d, q_d, r_d, s_d, u_d1, u_d2, &
121 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n) &
122 bind(c, name = 'hip_pipecg_vecops')
123 use, intrinsic :: iso_c_binding
124 import c_rp
125 implicit none
126 type(c_ptr), value :: p_d, q_d, r_d, s_d, u_d1, u_d2
127 type(c_ptr), value :: w_d, ni_d, mi_d, z_d, mult_d
128 integer(c_int) :: n
129 real(c_rp) :: alpha, beta, reduction(3)
130 end subroutine hip_pipecg_vecops
131 end interface
132
133 interface
134 subroutine hip_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, &
135 p_cur, p_space, n) &
136 bind(c, name = 'hip_cg_update_xp')
137 use, intrinsic :: iso_c_binding
138 implicit none
139 type(c_ptr), value :: x_d, p_d, u_d_d, alpha, beta
140 integer(c_int) :: p_cur, n, p_space
141 end subroutine hip_cg_update_xp
142 end interface
143#endif
144
145contains
146
147 subroutine device_pipecg_vecops(p_d, q_d, r_d, s_d, u_d1, u_d2, &
148 w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction,n)
149 type(c_ptr), value :: p_d, q_d, r_d, s_d, u_d1, u_d2
150 type(c_ptr), value :: w_d, ni_d, mi_d, z_d, mult_d
151 integer(c_int) :: n
152 real(c_rp) :: alpha, beta, reduction(3)
153#ifdef HAVE_HIP
154 call hip_pipecg_vecops(p_d, q_d, r_d,&
155 s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, &
156 mult_d, reduction,n)
157#elif HAVE_CUDA
158 call cuda_pipecg_vecops(p_d, q_d, r_d,&
159 s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, &
160 mult_d, reduction,n)
161#else
162 call neko_error('No device backend configured')
163#endif
164 end subroutine device_pipecg_vecops
165
166 subroutine device_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, &
167 p_cur, p_space, n)
168 use, intrinsic :: iso_c_binding
169 type(c_ptr), value :: x_d, p_d, u_d_d, alpha, beta
170 integer(c_int) :: p_cur, n, p_space
171#ifdef HAVE_HIP
172 call hip_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, p_cur, p_space, n)
173#elif HAVE_CUDA
174 call cuda_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, p_cur, p_space, n)
175#else
176 call neko_error('No device backend configured')
177#endif
178 end subroutine device_cg_update_xp
179
181 subroutine pipecg_device_init(this, n, max_iter, M, rel_tol, abs_tol, &
182 monitor)
183 class(pipecg_device_t), target, intent(inout) :: this
184 class(pc_t), optional, intent(in), target :: M
185 integer, intent(in) :: n
186 integer, intent(in) :: max_iter
187 real(kind=rp), optional, intent(in) :: rel_tol
188 real(kind=rp), optional, intent(in) :: abs_tol
189 logical, optional, intent(in) :: monitor
190 type(c_ptr) :: ptr
191 integer(c_size_t) :: u_size
192 integer :: i
193
194 call this%free()
195
196 allocate(this%p(n))
197 allocate(this%q(n))
198 allocate(this%r(n))
199 allocate(this%s(n))
200 allocate(this%u(n, device_pipecg_p_space+1))
201 allocate(this%u_d(device_pipecg_p_space+1))
202 allocate(this%w(n))
203 allocate(this%z(n))
204 allocate(this%mi(n))
205 allocate(this%ni(n))
206 allocate(this%alpha(device_pipecg_p_space))
207 allocate(this%beta(device_pipecg_p_space))
208
209 if (present(m)) then
210 this%M => m
211 end if
212
213 call device_map(this%p, this%p_d, n)
214 call device_map(this%q, this%q_d, n)
215 call device_map(this%r, this%r_d, n)
216 call device_map(this%s, this%s_d, n)
217 call device_map(this%w, this%w_d, n)
218 call device_map(this%z, this%z_d, n)
219 call device_map(this%mi, this%mi_d, n)
220 call device_map(this%ni, this%ni_d, n)
221 call device_map(this%alpha, this%alpha_d, device_pipecg_p_space)
222 call device_map(this%beta, this%beta_d, device_pipecg_p_space)
223 do i = 1, device_pipecg_p_space+1
224 this%u_d(i) = c_null_ptr
225 call device_map(this%u(:,i), this%u_d(i), n)
226 end do
227 !Did not work with 4 for some reason...
228 u_size = 8*(device_pipecg_p_space+1)
229 call device_alloc(this%u_d_d, u_size)
230 ptr = c_loc(this%u_d)
231 call device_memcpy(ptr,this%u_d_d, u_size, &
232 host_to_device, sync = .false.)
233
234 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
235 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
236 else if (present(rel_tol) .and. present(abs_tol)) then
237 call this%ksp_init(max_iter, rel_tol, abs_tol)
238 else if (present(monitor) .and. present(abs_tol)) then
239 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
240 else if (present(rel_tol) .and. present(monitor)) then
241 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
242 else if (present(rel_tol)) then
243 call this%ksp_init(max_iter, rel_tol = rel_tol)
244 else if (present(abs_tol)) then
245 call this%ksp_init(max_iter, abs_tol = abs_tol)
246 else if (present(monitor)) then
247 call this%ksp_init(max_iter, monitor = monitor)
248 else
249 call this%ksp_init(max_iter)
250 end if
251
252 call device_event_create(this%gs_event, 2)
253
254 end subroutine pipecg_device_init
255
257 subroutine pipecg_device_free(this)
258 class(pipecg_device_t), intent(inout) :: this
259 integer :: i
260
261 call this%ksp_free()
262
263 if (allocated(this%p)) then
264 if (c_associated(this%p_d)) then
265 call device_unmap(this%p, this%p_d)
266 end if
267 deallocate(this%p)
268 end if
269 if (allocated(this%q)) then
270 if (c_associated(this%q_d)) then
271 call device_unmap(this%q, this%q_d)
272 end if
273 deallocate(this%q)
274 end if
275 if (allocated(this%r)) then
276 if (c_associated(this%r_d)) then
277 call device_unmap(this%r, this%r_d)
278 end if
279 deallocate(this%r)
280 end if
281 if (allocated(this%s)) then
282 if (c_associated(this%s_d)) then
283 call device_unmap(this%s, this%s_d)
284 end if
285 deallocate(this%s)
286 end if
287 if (allocated(this%u)) then
288 if (allocated(this%u_d)) then
289 do i = 1, device_pipecg_p_space+1
290 if (c_associated(this%u_d(i))) then
291 call device_unmap(this%u(:,i), this%u_d(i))
292 end if
293 end do
294 end if
295 deallocate(this%u)
296 end if
297 if (allocated(this%u_d)) then
298 deallocate(this%u_d)
299 end if
300 if (allocated(this%w)) then
301 if (c_associated(this%w_d)) then
302 call device_unmap(this%w, this%w_d)
303 end if
304 deallocate(this%w)
305 end if
306 if (allocated(this%z)) then
307 if (c_associated(this%z_d)) then
308 call device_unmap(this%z, this%z_d)
309 end if
310 deallocate(this%z)
311 end if
312 if (allocated(this%mi)) then
313 if (c_associated(this%mi_d)) then
314 call device_unmap(this%mi, this%mi_d)
315 end if
316 deallocate(this%mi)
317 end if
318 if (allocated(this%ni)) then
319 if (c_associated(this%ni_d)) then
320 call device_unmap(this%ni, this%ni_d)
321 end if
322 deallocate(this%ni)
323 end if
324 if (allocated(this%alpha)) then
325 if (c_associated(this%alpha_d)) then
326 call device_unmap(this%alpha, this%alpha_d)
327 end if
328 deallocate(this%alpha)
329 end if
330 if (allocated(this%beta)) then
331 if (c_associated(this%beta_d)) then
332 call device_unmap(this%beta, this%beta_d)
333 end if
334 deallocate(this%beta)
335 end if
336
337 if (c_associated(this%u_d_d)) then
338 call device_free(this%u_d_d)
339 end if
340
341 nullify(this%M)
342
343 if (c_associated(this%gs_event)) then
344 call device_event_destroy(this%gs_event)
345 end if
346
347 end subroutine pipecg_device_free
348
350 function pipecg_device_solve(this, Ax, x, f, n, coef, blst, gs_h, niter) result(ksp_results)
351 class(pipecg_device_t), intent(inout) :: this
352 class(ax_t), intent(in) :: ax
353 type(field_t), intent(inout) :: x
354 integer, intent(in) :: n
355 real(kind=rp), dimension(n), intent(in) :: f
356 type(coef_t), intent(inout) :: coef
357 type(bc_list_t), intent(inout) :: blst
358 type(gs_t), intent(inout) :: gs_h
359 type(ksp_monitor_t) :: ksp_results
360 integer, optional, intent(in) :: niter
361 integer :: iter, max_iter, ierr, p_cur, p_prev, u_prev
362 real(kind=rp) :: rnorm, rtr, reduction(3), norm_fac
363 real(kind=rp) :: gamma1, gamma2, delta
364 real(kind=rp) :: tmp1, tmp2, tmp3
365 type(mpi_request) :: request
366 type(mpi_status) :: status
367 type(c_ptr) :: f_d
368 f_d = device_get_ptr(f)
369
370 if (present(niter)) then
371 max_iter = niter
372 else
373 max_iter = this%max_iter
374 end if
375 norm_fac = 1.0_rp / sqrt(coef%volume)
376
377 associate(p => this%p, q => this%q, r => this%r, s => this%s, &
378 u => this%u, w => this%w, z => this%z, mi => this%mi, ni => this%ni, &
379 alpha => this%alpha, beta => this%beta, &
380 alpha_d => this%alpha_d, beta_d => this%beta_d, &
381 p_d => this%p_d, q_d => this%q_d, r_d => this%r_d, &
382 s_d => this%s_d, u_d => this%u_d, u_d_d => this%u_d_d, &
383 w_d => this%w_d, z_d => this%z_d, mi_d => this%mi_d, ni_d => this%ni_d)
384
385 p_prev = device_pipecg_p_space !this%p_space
386 u_prev = device_pipecg_p_space + 1 !this%p_space+1
387 p_cur = 1
388 call device_rzero(x%x_d, n)
389 call device_rzero(z_d, n)
390 call device_rzero(q_d, n)
391 call device_rzero(p_d, n)
392 call device_rzero(s_d, n)
393 call device_copy(r_d, f_d, n)
394 !apply u=M^-1r
395 !call device_copy(u_d(u_prev), r_d, n)
396 call this%M%solve(u(1, u_prev), r, n)
397 call ax%compute(w, u(1, u_prev), coef, x%msh, x%Xh)
398 call gs_h%op(w, n, gs_op_add, this%gs_event)
399 call device_event_sync(this%gs_event)
400 call blst%apply_scalar(w, n)
401
402 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
403 rnorm = sqrt(rtr)*norm_fac
404 ksp_results%res_start = rnorm
405 ksp_results%res_final = rnorm
406 ksp_results%iter = 0
407 if (abscmp(rnorm, 0.0_rp)) then
408 ksp_results%converged = .true.
409 return
410 end if
411
412 gamma1 = 0.0_rp
413 tmp1 = 0.0_rp
414 tmp2 = 0.0_rp
415 tmp3 = 0.0_rp
416 tmp1 = device_vlsc3(r_d, coef%mult_d, u_d(u_prev), n)
417 tmp2 = device_vlsc3(w_d, coef%mult_d, u_d(u_prev), n)
418 tmp3 = device_vlsc3(r_d, coef%mult_d, r_d, n)
419 reduction(1) = tmp1
420 reduction(2) = tmp2
421 reduction(3) = tmp3
422
423 call this%monitor_start('PipeCG')
424 do iter = 1, max_iter
425 call mpi_iallreduce(mpi_in_place, reduction, 3, &
426 mpi_real_precision, mpi_sum, neko_comm, request, ierr)
427
428 call this%M%solve(mi, w, n)
429 call ax%compute(ni, mi, coef, x%msh, x%Xh)
430 call gs_h%op(ni, n, gs_op_add, this%gs_event)
431 call device_event_sync(this%gs_event)
432 call blst%apply(ni, n)
433
434 call mpi_wait(request, status, ierr)
435 gamma2 = gamma1
436 gamma1 = reduction(1)
437 delta = reduction(2)
438 rtr = reduction(3)
439
440 rnorm = sqrt(rtr)*norm_fac
441 call this%monitor_iter(iter, rnorm)
442 if (rnorm .lt. this%abs_tol) exit
443
444
445 if (iter .gt. 1) then
446 beta(p_cur) = gamma1 / gamma2
447 alpha(p_cur) = gamma1 / (delta - (beta(p_cur) * gamma1/alpha(p_prev)))
448 else
449 beta(p_cur) = 0.0_rp
450 alpha(p_cur) = gamma1/delta
451 end if
452
453 call device_pipecg_vecops(p_d, q_d, r_d,&
454 s_d, u_d(u_prev), u_d(p_cur),&
455 w_d, z_d, ni_d,&
456 mi_d, alpha(p_cur), beta(p_cur),&
457 coef%mult_d, reduction, n)
458 if (p_cur .eq. device_pipecg_p_space) then
459 call device_memcpy(alpha, alpha_d, p_cur, &
460 host_to_device, sync=.false.)
461 call device_memcpy(beta, beta_d, p_cur, &
462 host_to_device, sync=.false.)
463 call device_cg_update_xp(x%x_d, p_d, u_d_d, alpha_d, beta_d, p_cur, &
465 p_prev = p_cur
466 u_prev = device_pipecg_p_space + 1
467 alpha(1) = alpha(p_cur)
468 beta(1) = beta(p_cur)
469 p_cur = 1
470 else
471 u_prev = p_cur
472 p_prev = p_cur
473 p_cur = p_cur + 1
474 end if
475 end do
476
477 if ( p_cur .ne. 1) then
478 call device_memcpy(alpha, alpha_d, p_cur, host_to_device, sync=.false.)
479 call device_memcpy(beta, beta_d, p_cur, host_to_device, sync=.false.)
480 call device_cg_update_xp(x%x_d, p_d, u_d_d, alpha_d, beta_d, p_cur, &
482 end if
483 call this%monitor_stop()
484 ksp_results%res_final = rnorm
485 ksp_results%iter = iter
486 ksp_results%converged = this%is_converged(iter, rnorm)
487
488 end associate
489
490 end function pipecg_device_solve
491
493 function pipecg_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
494 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
495 class(pipecg_device_t), intent(inout) :: this
496 class(ax_t), intent(in) :: ax
497 type(field_t), intent(inout) :: x
498 type(field_t), intent(inout) :: y
499 type(field_t), intent(inout) :: z
500 integer, intent(in) :: n
501 real(kind=rp), dimension(n), intent(in) :: fx
502 real(kind=rp), dimension(n), intent(in) :: fy
503 real(kind=rp), dimension(n), intent(in) :: fz
504 type(coef_t), intent(inout) :: coef
505 type(bc_list_t), intent(inout) :: blstx
506 type(bc_list_t), intent(inout) :: blsty
507 type(bc_list_t), intent(inout) :: blstz
508 type(gs_t), intent(inout) :: gs_h
509 type(ksp_monitor_t), dimension(3) :: ksp_results
510 integer, optional, intent(in) :: niter
511
512 ksp_results(1) = this%solve(ax, x, fx, n, coef, blstx, gs_h, niter)
513 ksp_results(2) = this%solve(ax, y, fy, n, coef, blsty, gs_h, niter)
514 ksp_results(3) = this%solve(ax, z, fz, n, coef, blstz, gs_h, niter)
515
516 end function pipecg_device_solve_coupled
517
518end module pipecg_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
Coefficients.
Definition coef.f90:34
Definition comm.F90:1
type(mpi_datatype), public mpi_real_precision
MPI type for working precision of REAL types.
Definition comm.F90:54
integer, public pe_size
MPI size of communicator.
Definition comm.F90:62
type(mpi_comm), public neko_comm
MPI communicator.
Definition comm.F90:46
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_vlsc3(u_d, v_d, w_d, n, strm)
Compute multiplication sum .
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
integer, parameter, public host_to_device
Definition device.F90:48
subroutine, public device_free(x_d)
Deallocate memory on the device.
Definition device.F90:243
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1623
subroutine, public device_alloc(x_d, s)
Allocate memory on the device.
Definition device.F90:212
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
integer, parameter, public ksp_max_iter
Maximum number of iters.
Definition krylov.f90:51
Definition math.f90:60
real(kind=rp) function, public glsc3(a, b, c, n)
Weighted inner product .
Definition math.f90:1287
subroutine, public copy(a, b, n)
Copy a vector .
Definition math.f90:291
subroutine, public rzero(a, n)
Zero a real vector.
Definition math.f90:235
integer, parameter, public c_rp
Definition num_types.f90:13
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Defines a pipelined Conjugate Gradient methods.
subroutine device_pipecg_vecops(p_d, q_d, r_d, s_d, u_d1, u_d2, w_d, z_d, ni_d, mi_d, alpha, beta, mult_d, reduction, n)
subroutine pipecg_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a pipelined PCG solver.
subroutine device_cg_update_xp(x_d, p_d, u_d_d, alpha, beta, p_cur, p_space, n)
integer, parameter device_pipecg_p_space
type(ksp_monitor_t) function, dimension(3) pipecg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, blstx, blsty, blstz, gs_h, niter)
Pipelined PCG coupled solve.
subroutine pipecg_device_free(this)
Deallocate a pipelined PCG solver.
type(ksp_monitor_t) function pipecg_device_solve(this, ax, x, f, n, coef, blst, gs_h, niter)
Pipelined PCG solve.
Krylov preconditioner.
Definition precon.f90:34
Utilities.
Definition utils.f90:35
void hip_cg_update_xp(void *x, void *p, void *u, void *alpha, void *beta, int *p_cur, int *p_space, int *n)
void hip_pipecg_vecops(void *p, void *q, void *r, void *s, void *u1, void *u2, void *w, void *z, void *ni, void *mi, real *alpha, real *beta, void *mult, real *reduction, int *n)
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
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
Pipelined preconditioned conjugate gradient method.
Defines a canonical Krylov preconditioner.
Definition precon.f90:40