57 use mpi_f08,
only : mpi_in_place, mpi_sum, mpi_allreduce
58 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_loc, &
59 c_associated, c_int, c_size_t, c_sizeof
65 integer :: m_restart = 30
66 real(kind=
rp),
allocatable :: w(:)
67 real(kind=
rp),
allocatable :: c(:)
68 real(kind=
rp),
allocatable :: r(:)
69 real(kind=
rp),
allocatable :: z(:,:)
70 real(kind=
rp),
allocatable :: h(:,:)
71 real(kind=
rp),
allocatable :: v(:,:)
72 real(kind=
rp),
allocatable :: s(:)
73 real(kind=
rp),
allocatable :: gam(:)
74 type(c_ptr) :: w_d = c_null_ptr
75 type(c_ptr) :: c_d = c_null_ptr
76 type(c_ptr) :: r_d = c_null_ptr
77 type(c_ptr) :: s_d = c_null_ptr
78 type(c_ptr) :: gam_d = c_null_ptr
79 type(c_ptr),
allocatable :: z_d(:), h_d(:), v_d(:)
80 type(c_ptr) :: z_d_d = c_null_ptr
81 type(c_ptr) :: h_d_d = c_null_ptr
82 type(c_ptr) :: v_d_d = c_null_ptr
83 type(c_ptr) :: gs_event = c_null_ptr
94 bind(c, name =
'hip_gmres_part2')
95 use,
intrinsic :: iso_c_binding
98 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
99 integer(c_int) :: j, n
106 bind(c, name =
'cuda_gmres_part2')
107 use,
intrinsic :: iso_c_binding
110 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
111 integer(c_int) :: j, n
119 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
120 integer(c_int) :: j, n
128 call neko_error(
'No device backend configured')
131#ifndef HAVE_DEVICE_MPI
132 if (pe_size .gt. 1)
then
133 call mpi_allreduce(mpi_in_place, alpha, 1, &
134 mpi_real_precision, mpi_sum, neko_comm, ierr)
143 integer,
intent(in) :: n
144 integer,
intent(in) :: max_iter
145 class(pc_t),
optional,
intent(in),
target :: M
146 real(kind=rp),
optional,
intent(in) :: rel_tol
147 real(kind=rp),
optional,
intent(in) :: abs_tol
148 logical,
optional,
intent(in) :: monitor
149 type(device_ident_t),
target :: M_ident
151 integer(c_size_t) :: z_size
164 call device_map(this%w, this%w_d, n)
165 call device_map(this%r, this%r_d, n)
167 allocate(this%c(this%m_restart))
168 allocate(this%s(this%m_restart))
169 allocate(this%gam(this%m_restart + 1))
170 call device_map(this%c, this%c_d, this%m_restart)
171 call device_map(this%s, this%s_d, this%m_restart)
172 call device_map(this%gam, this%gam_d, this%m_restart+1)
174 allocate(this%z(n, this%m_restart))
175 allocate(this%v(n, this%m_restart))
176 allocate(this%h(this%m_restart, this%m_restart))
177 allocate(this%z_d(this%m_restart))
178 allocate(this%v_d(this%m_restart))
179 allocate(this%h_d(this%m_restart))
180 do i = 1, this%m_restart
181 this%z_d(i) = c_null_ptr
182 call device_map(this%z(:,i), this%z_d(i), n)
184 this%v_d(i) = c_null_ptr
185 call device_map(this%v(:,i), this%v_d(i), n)
187 this%h_d(i) = c_null_ptr
188 call device_map(this%h(:,i), this%h_d(i), this%m_restart)
191 z_size = c_sizeof(c_null_ptr) * (this%m_restart)
192 call device_alloc(this%z_d_d, z_size)
193 call device_alloc(this%v_d_d, z_size)
194 call device_alloc(this%h_d_d, z_size)
195 ptr = c_loc(this%z_d)
196 call device_memcpy(ptr, this%z_d_d, z_size, &
197 host_to_device, sync = .false.)
198 ptr = c_loc(this%v_d)
199 call device_memcpy(ptr, this%v_d_d, z_size, &
200 host_to_device, sync = .false.)
201 ptr = c_loc(this%h_d)
202 call device_memcpy(ptr, this%h_d_d, z_size, &
203 host_to_device, sync = .false.)
206 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
207 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
208 else if (
present(rel_tol) .and.
present(abs_tol))
then
209 call this%ksp_init(max_iter, rel_tol, abs_tol)
210 else if (
present(monitor) .and.
present(abs_tol))
then
211 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
212 else if (
present(rel_tol) .and.
present(monitor))
then
213 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
214 else if (
present(rel_tol))
then
215 call this%ksp_init(max_iter, rel_tol = rel_tol)
216 else if (
present(abs_tol))
then
217 call this%ksp_init(max_iter, abs_tol = abs_tol)
218 else if (
present(monitor))
then
219 call this%ksp_init(max_iter, monitor = monitor)
221 call this%ksp_init(max_iter)
224 call device_event_create(this%gs_event, 2)
235 if (
allocated(this%w))
then
236 if (c_associated(this%w_d))
then
237 call device_unmap(this%w, this%w_d)
242 if (
allocated(this%c))
then
243 if (c_associated(this%c_d))
then
244 call device_unmap(this%c, this%c_d)
249 if (
allocated(this%r))
then
250 if (c_associated(this%r_d))
then
251 call device_unmap(this%r, this%r_d)
256 if (
allocated(this%z))
then
257 if (
allocated(this%z_d))
then
258 do i = 1, this%m_restart
259 if (c_associated(this%z_d(i)))
then
260 call device_unmap(this%z(:,i), this%z_d(i))
267 if (
allocated(this%h))
then
268 if (
allocated(this%h_d))
then
269 do i = 1, this%m_restart
270 if (c_associated(this%h_d(i)))
then
271 call device_unmap(this%h(:,i), this%h_d(i))
278 if (
allocated(this%v))
then
279 if (
allocated(this%v_d))
then
280 do i = 1, this%m_restart
281 if (c_associated(this%v_d(i)))
then
282 call device_unmap(this%v(:,i), this%v_d(i))
289 if (
allocated(this%s))
then
290 if (c_associated(this%s_d))
then
291 call device_unmap(this%s, this%s_d)
295 if (
allocated(this%gam))
then
296 if (c_associated(this%gam_d))
then
297 call device_unmap(this%gam, this%gam_d)
302 if (c_associated(this%z_d_d))
then
303 call device_free(this%z_d_d)
305 if (c_associated(this%v_d_d))
then
306 call device_free(this%v_d_d)
308 if (c_associated(this%h_d_d))
then
309 call device_free(this%h_d_d)
314 if (c_associated(this%gs_event))
then
315 call device_event_destroy(this%gs_event)
322 niter)
result(ksp_results)
324 class(ax_t),
intent(in) :: ax
325 type(field_t),
intent(inout) :: x
326 integer,
intent(in) :: n
327 real(kind=rp),
dimension(n),
intent(in) :: f
328 type(coef_t),
intent(inout) :: coef
329 class(scalar_bc_projector_t),
intent(inout) :: bc_projector
330 type(gs_t),
intent(inout) :: gs_h
331 type(ksp_monitor_t) :: ksp_results
332 integer,
optional,
intent(in) :: niter
333 integer :: iter, max_iter
335 real(kind=rp) :: rnorm, alpha, temp, lr, alpha2, norm_fac
339 f_d = device_get_ptr(f)
345 if (
present(niter))
then
348 max_iter = this%max_iter
351 associate(w => this%w, c => this%c, r => this%r, z => this%z, h => this%h, &
352 v => this%v, s => this%s, gam => this%gam, v_d => this%v_d, &
353 w_d => this%w_d, r_d => this%r_d, h_d => this%h_d, &
354 v_d_d => this%v_d_d, x_d => x%x_d, z_d_d => this%z_d_d, &
357 norm_fac = 1.0_rp / sqrt(coef%volume)
358 call rzero(gam, this%m_restart + 1)
359 call rone(s, this%m_restart)
360 call rone(c, this%m_restart)
361 call rzero(h, this%m_restart * this%m_restart)
362 call device_rzero(x%x_d, n)
363 call device_rzero(this%gam_d, this%m_restart + 1)
364 call device_rone(this%s_d, this%m_restart)
365 call device_rone(this%c_d, this%m_restart)
367 call rzero(this%h, this%m_restart**2)
372 call this%monitor_start(
'GMRES')
373 do while (.not. conv .and. iter .lt. max_iter)
375 if (iter .eq. 0)
then
376 call device_copy(r_d, f_d, n)
378 call device_copy(r_d, f_d, n)
379 call ax%compute(w, x%x, coef, x%msh, x%Xh)
380 call gs_h%op(w, n, gs_op_add, this%gs_event)
381 call device_event_sync(this%gs_event)
382 call bc_projector%apply(w, n)
383 call device_sub2(r_d, w_d, n)
386 gam(1) = sqrt(device_glsc3(r_d, r_d, coef%mult_d, n))
387 if (iter .eq. 0)
then
388 ksp_results%res_start = gam(1) * norm_fac
391 if (abscmp(gam(1), 0.0_rp))
exit
394 temp = 1.0_rp / gam(1)
395 call device_cmult2(v_d(1), r_d, temp, n)
396 do j = 1, this%m_restart
399 call this%M%solve(z(1,j), v(1,j), n)
401 call ax%compute(w, z(1,j), coef, x%msh, x%Xh)
402 call gs_h%op(w, n, gs_op_add, this%gs_event)
403 call device_event_sync(this%gs_event)
404 call bc_projector%apply(w, n)
406 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
408 h(i,j) = device_glsc3(w_d, v_d(i), coef%mult_d, n)
410 call device_add2s2(w_d, v_d(i), -h(i,j), n)
412 alpha2 = device_glsc3(w_d, w_d, coef%mult_d, n)
415 call device_glsc3_many(h(1,j), w_d, v_d_d, coef%mult_d, j, n)
417 call device_memcpy(h(:,j), h_d(j), j, &
418 host_to_device, sync = .false.)
428 h(i,j) = c(i)*temp + s(i) * h(i+1,j)
429 h(i+1,j) = -s(i)*temp + c(i) * h(i+1,j)
433 if (abscmp(alpha, 0.0_rp))
then
438 lr = sqrt(h(j,j) * h(j,j) + alpha2)
443 call device_memcpy(h(:,j), h_d(j), j, &
444 host_to_device, sync = .false.)
445 gam(j+1) = -s(j) * gam(j)
446 gam(j) = c(j) * gam(j)
448 rnorm = abs(gam(j+1)) * norm_fac
449 call this%monitor_iter(iter, rnorm)
450 if (rnorm .lt. this%abs_tol)
then
455 if (iter + 1 .gt. max_iter)
exit
457 if (j .lt. this%m_restart)
then
458 temp = 1.0_rp / alpha
459 call device_cmult2(v_d(j+1), w_d, temp, n)
464 j = min(j, this%m_restart)
468 temp = temp - h(k,i) * c(i)
473 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
475 call device_add2s2(x_d, this%z_d(i), c(i), n)
478 call device_memcpy(c, c_d, j, host_to_device, sync = .false.)
479 call device_add2s2_many(x_d, z_d_d, c_d, j, n)
484 call this%monitor_stop()
485 ksp_results%res_final = rnorm
486 ksp_results%iter = iter
487 ksp_results%converged = this%is_converged(iter, rnorm)
493 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
495 class(ax_t),
intent(in) :: ax
496 type(field_t),
intent(inout) :: x
497 type(field_t),
intent(inout) :: y
498 type(field_t),
intent(inout) :: z
499 integer,
intent(in) :: n
500 real(kind=rp),
dimension(n),
intent(in) :: fx
501 real(kind=rp),
dimension(n),
intent(in) :: fy
502 real(kind=rp),
dimension(n),
intent(in) :: fz
503 type(coef_t),
intent(inout) :: coef
504 class(vector_bc_projector_t),
intent(inout) :: bc_projector
505 type(gs_t),
intent(inout) :: gs_h
506 type(ksp_monitor_t),
dimension(3) :: ksp_results
507 integer,
optional,
intent(in) :: niter
508 type(scalar_bc_projector_t),
pointer :: bc_x, bc_y, bc_z
510 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
511 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
512 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
513 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
real cuda_gmres_part2(void *w, void *v, void *h, void *mult, int *j, int *n)
Return the device pointer for an associated Fortran array.
Map a Fortran array to a device (allocate and associate)
Copy data between host and device (or device and device)
Unmap a Fortran array from a device (deassociate and free)
Defines a Matrix-vector product.
type(mpi_datatype), public mpi_real_precision
MPI type for working precision of REAL types.
integer, public pe_size
MPI size of communicator.
type(mpi_comm), public neko_comm
MPI communicator.
Identity Krylov preconditioner for accelerators.
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_add2s2_many(y_d, x_d_d, a_d, j, n, strm)
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_rone(a_d, n, strm)
Set all elements to one.
subroutine, public device_glsc3_many(h, w_d, v_d_d, mult_d, j, n, strm)
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.
subroutine, public device_event_sync(event)
Synchronize an event.
integer, parameter, public host_to_device
subroutine, public device_free(x_d)
Deallocate memory on the device.
subroutine, public device_event_destroy(event)
Destroy a device event.
subroutine, public device_alloc(x_d, s)
Allocate memory on the device.
subroutine, public device_event_create(event, flags)
Create a device event queue.
Defines various GMRES methods.
real(c_rp) function device_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
subroutine gmres_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a standard GMRES solver.
type(ksp_monitor_t) function, dimension(3) gmres_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard GMRES coupled solve.
type(ksp_monitor_t) function gmres_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Standard GMRES solve.
subroutine gmres_device_free(this)
Deallocate a standard GMRES solver.
Implements the base abstract type for Krylov solvers plus helper types.
subroutine, public rone(a, n)
Set all elements to one.
subroutine, public rzero(a, n)
Zero a real vector.
integer, parameter neko_bcknd_opencl
integer, parameter neko_bcknd_metal
integer, parameter, public c_rp
integer, parameter, public rp
Global precision used in computations.
Implements scalar_projector_t.
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 .
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Defines a canonical Krylov preconditioner for accelerators.
Standard preconditioned generalized minimal residual method.
Type for storing initial and final residuals in a Krylov solver.
Base abstract type for a canonical Krylov method, solving .
Defines a canonical Krylov preconditioner.
Projector for scalar boundary conditions.
Abstract type for resolving vector boundary conditions.