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