55 use mpi_f08,
only : mpi_in_place, mpi_sum, mpi_allreduce
56 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_loc, &
57 c_associated, c_int, c_size_t, c_sizeof
63 integer :: m_restart = 30
64 real(kind=
rp),
allocatable :: w(:)
65 real(kind=
rp),
allocatable :: c(:)
66 real(kind=
rp),
allocatable :: r(:)
67 real(kind=
rp),
allocatable :: z(:,:)
68 real(kind=
rp),
allocatable :: h(:,:)
69 real(kind=
rp),
allocatable :: v(:,:)
70 real(kind=
rp),
allocatable :: s(:)
71 real(kind=
rp),
allocatable :: gam(:)
72 type(c_ptr) :: w_d = c_null_ptr
73 type(c_ptr) :: c_d = c_null_ptr
74 type(c_ptr) :: r_d = c_null_ptr
75 type(c_ptr) :: s_d = c_null_ptr
76 type(c_ptr) :: gam_d = c_null_ptr
77 type(c_ptr),
allocatable :: z_d(:), h_d(:), v_d(:)
78 type(c_ptr) :: z_d_d = c_null_ptr
79 type(c_ptr) :: h_d_d = c_null_ptr
80 type(c_ptr) :: v_d_d = c_null_ptr
81 type(c_ptr) :: gs_event = c_null_ptr
92 bind(c, name =
'hip_gmres_part2')
93 use,
intrinsic :: iso_c_binding
96 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
97 integer(c_int) :: j, n
104 bind(c, name =
'cuda_gmres_part2')
105 use,
intrinsic :: iso_c_binding
108 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
109 integer(c_int) :: j, n
117 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
118 integer(c_int) :: j, n
126 call neko_error(
'No device backend configured')
129#ifndef HAVE_DEVICE_MPI
130 if (pe_size .gt. 1)
then
131 call mpi_allreduce(mpi_in_place, alpha, 1, &
132 mpi_real_precision, mpi_sum, neko_comm, ierr)
141 integer,
intent(in) :: n
142 integer,
intent(in) :: max_iter
143 class(pc_t),
optional,
intent(in),
target :: M
144 real(kind=rp),
optional,
intent(in) :: rel_tol
145 real(kind=rp),
optional,
intent(in) :: abs_tol
146 logical,
optional,
intent(in) :: monitor
147 type(device_ident_t),
target :: M_ident
149 integer(c_size_t) :: z_size
162 call device_map(this%w, this%w_d, n)
163 call device_map(this%r, this%r_d, n)
165 allocate(this%c(this%m_restart))
166 allocate(this%s(this%m_restart))
167 allocate(this%gam(this%m_restart + 1))
168 call device_map(this%c, this%c_d, this%m_restart)
169 call device_map(this%s, this%s_d, this%m_restart)
170 call device_map(this%gam, this%gam_d, this%m_restart+1)
172 allocate(this%z(n, this%m_restart))
173 allocate(this%v(n, this%m_restart))
174 allocate(this%h(this%m_restart, this%m_restart))
175 allocate(this%z_d(this%m_restart))
176 allocate(this%v_d(this%m_restart))
177 allocate(this%h_d(this%m_restart))
178 do i = 1, this%m_restart
179 this%z_d(i) = c_null_ptr
180 call device_map(this%z(:,i), this%z_d(i), n)
182 this%v_d(i) = c_null_ptr
183 call device_map(this%v(:,i), this%v_d(i), n)
185 this%h_d(i) = c_null_ptr
186 call device_map(this%h(:,i), this%h_d(i), this%m_restart)
189 z_size = c_sizeof(c_null_ptr) * (this%m_restart)
190 call device_alloc(this%z_d_d, z_size)
191 call device_alloc(this%v_d_d, z_size)
192 call device_alloc(this%h_d_d, z_size)
193 ptr = c_loc(this%z_d)
194 call device_memcpy(ptr, this%z_d_d, z_size, &
195 host_to_device, sync = .false.)
196 ptr = c_loc(this%v_d)
197 call device_memcpy(ptr, this%v_d_d, z_size, &
198 host_to_device, sync = .false.)
199 ptr = c_loc(this%h_d)
200 call device_memcpy(ptr, this%h_d_d, z_size, &
201 host_to_device, sync = .false.)
204 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
205 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
206 else if (
present(rel_tol) .and.
present(abs_tol))
then
207 call this%ksp_init(max_iter, rel_tol, abs_tol)
208 else if (
present(monitor) .and.
present(abs_tol))
then
209 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
210 else if (
present(rel_tol) .and.
present(monitor))
then
211 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
212 else if (
present(rel_tol))
then
213 call this%ksp_init(max_iter, rel_tol = rel_tol)
214 else if (
present(abs_tol))
then
215 call this%ksp_init(max_iter, abs_tol = abs_tol)
216 else if (
present(monitor))
then
217 call this%ksp_init(max_iter, monitor = monitor)
219 call this%ksp_init(max_iter)
222 call device_event_create(this%gs_event, 2)
233 if (
allocated(this%w))
then
234 if (c_associated(this%w_d))
then
235 call device_unmap(this%w, this%w_d)
240 if (
allocated(this%c))
then
241 if (c_associated(this%c_d))
then
242 call device_unmap(this%c, this%c_d)
247 if (
allocated(this%r))
then
248 if (c_associated(this%r_d))
then
249 call device_unmap(this%r, this%r_d)
254 if (
allocated(this%z))
then
255 if (
allocated(this%z_d))
then
256 do i = 1, this%m_restart
257 if (c_associated(this%z_d(i)))
then
258 call device_unmap(this%z(:,i), this%z_d(i))
265 if (
allocated(this%h))
then
266 if (
allocated(this%h_d))
then
267 do i = 1, this%m_restart
268 if (c_associated(this%h_d(i)))
then
269 call device_unmap(this%h(:,i), this%h_d(i))
276 if (
allocated(this%v))
then
277 if (
allocated(this%v_d))
then
278 do i = 1, this%m_restart
279 if (c_associated(this%v_d(i)))
then
280 call device_unmap(this%v(:,i), this%v_d(i))
287 if (
allocated(this%s))
then
288 if (c_associated(this%s_d))
then
289 call device_unmap(this%s, this%s_d)
293 if (
allocated(this%gam))
then
294 if (c_associated(this%gam_d))
then
295 call device_unmap(this%gam, this%gam_d)
300 if (c_associated(this%z_d_d))
then
301 call device_free(this%z_d_d)
303 if (c_associated(this%v_d_d))
then
304 call device_free(this%v_d_d)
306 if (c_associated(this%h_d_d))
then
307 call device_free(this%h_d_d)
312 if (c_associated(this%gs_event))
then
313 call device_event_destroy(this%gs_event)
322 class(ax_t),
intent(in) :: ax
323 type(field_t),
intent(inout) :: x
324 integer,
intent(in) :: n
325 real(kind=rp),
dimension(n),
intent(in) :: f
326 type(coef_t),
intent(inout) :: coef
327 type(bc_list_t),
intent(inout) :: blst
328 type(gs_t),
intent(inout) :: gs_h
329 type(ksp_monitor_t) :: ksp_results
330 integer,
optional,
intent(in) :: niter
331 integer :: iter, max_iter
333 real(kind=rp) :: rnorm, alpha, temp, lr, alpha2, norm_fac
337 f_d = device_get_ptr(f)
343 if (
present(niter))
then
346 max_iter = this%max_iter
349 associate(w => this%w, c => this%c, r => this%r, z => this%z, h => this%h, &
350 v => this%v, s => this%s, gam => this%gam, v_d => this%v_d, &
351 w_d => this%w_d, r_d => this%r_d, h_d => this%h_d, &
352 v_d_d => this%v_d_d, x_d => x%x_d, z_d_d => this%z_d_d, &
355 norm_fac = 1.0_rp / sqrt(coef%volume)
356 call rzero(gam, this%m_restart + 1)
357 call rone(s, this%m_restart)
358 call rone(c, this%m_restart)
359 call rzero(h, this%m_restart * this%m_restart)
360 call device_rzero(x%x_d, n)
361 call device_rzero(this%gam_d, this%m_restart + 1)
362 call device_rone(this%s_d, this%m_restart)
363 call device_rone(this%c_d, this%m_restart)
365 call rzero(this%h, this%m_restart**2)
370 call this%monitor_start(
'GMRES')
371 do while (.not. conv .and. iter .lt. max_iter)
373 if (iter .eq. 0)
then
374 call device_copy(r_d, f_d, n)
376 call device_copy(r_d, f_d, n)
377 call ax%compute(w, x%x, coef, x%msh, x%Xh)
378 call gs_h%op(w, n, gs_op_add, this%gs_event)
379 call device_event_sync(this%gs_event)
380 call blst%apply_scalar(w, n)
381 call device_sub2(r_d, w_d, n)
384 gam(1) = sqrt(device_glsc3(r_d, r_d, coef%mult_d, n))
385 if (iter .eq. 0)
then
386 ksp_results%res_start = gam(1) * norm_fac
389 if (abscmp(gam(1), 0.0_rp))
exit
392 temp = 1.0_rp / gam(1)
393 call device_cmult2(v_d(1), r_d, temp, n)
394 do j = 1, this%m_restart
397 call this%M%solve(z(1,j), v(1,j), n)
399 call ax%compute(w, z(1,j), coef, x%msh, x%Xh)
400 call gs_h%op(w, n, gs_op_add, this%gs_event)
401 call device_event_sync(this%gs_event)
402 call blst%apply_scalar(w, n)
404 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
406 h(i,j) = device_glsc3(w_d, v_d(i), coef%mult_d, n)
408 call device_add2s2(w_d, v_d(i), -h(i,j), n)
410 alpha2 = device_glsc3(w_d, w_d, coef%mult_d, n)
413 call device_glsc3_many(h(1,j), w_d, v_d_d, coef%mult_d, j, n)
415 call device_memcpy(h(:,j), h_d(j), j, &
416 host_to_device, sync = .false.)
426 h(i,j) = c(i)*temp + s(i) * h(i+1,j)
427 h(i+1,j) = -s(i)*temp + c(i) * h(i+1,j)
431 if (abscmp(alpha, 0.0_rp))
then
436 lr = sqrt(h(j,j) * h(j,j) + alpha2)
441 call device_memcpy(h(:,j), h_d(j), j, &
442 host_to_device, sync = .false.)
443 gam(j+1) = -s(j) * gam(j)
444 gam(j) = c(j) * gam(j)
446 rnorm = abs(gam(j+1)) * norm_fac
447 call this%monitor_iter(iter, rnorm)
448 if (rnorm .lt. this%abs_tol)
then
453 if (iter + 1 .gt. max_iter)
exit
455 if (j .lt. this%m_restart)
then
456 temp = 1.0_rp / alpha
457 call device_cmult2(v_d(j+1), w_d, temp, n)
462 j = min(j, this%m_restart)
466 temp = temp - h(k,i) * c(i)
471 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
473 call device_add2s2(x_d, this%z_d(i), c(i), n)
476 call device_memcpy(c, c_d, j, host_to_device, sync = .false.)
477 call device_add2s2_many(x_d, z_d_d, c_d, j, n)
482 call this%monitor_stop()
483 ksp_results%res_final = rnorm
484 ksp_results%iter = iter
485 ksp_results%converged = this%is_converged(iter, rnorm)
491 n, coef, blstx, blsty, blstz, gs_h, niter)
result(ksp_results)
493 class(ax_t),
intent(in) :: ax
494 type(field_t),
intent(inout) :: x
495 type(field_t),
intent(inout) :: y
496 type(field_t),
intent(inout) :: z
497 integer,
intent(in) :: n
498 real(kind=rp),
dimension(n),
intent(in) :: fx
499 real(kind=rp),
dimension(n),
intent(in) :: fy
500 real(kind=rp),
dimension(n),
intent(in) :: fz
501 type(coef_t),
intent(inout) :: coef
502 type(bc_list_t),
intent(inout) :: blstx
503 type(bc_list_t),
intent(inout) :: blsty
504 type(bc_list_t),
intent(inout) :: blstz
505 type(gs_t),
intent(inout) :: gs_h
506 type(ksp_monitor_t),
dimension(3) :: ksp_results
507 integer,
optional,
intent(in) :: niter
509 ksp_results(1) = this%solve(ax, x, fx, n, coef, blstx, gs_h, niter)
510 ksp_results(2) = this%solve(ax, y, fy, n, coef, blsty, gs_h, niter)
511 ksp_results(3) = this%solve(ax, z, fz, n, coef, blstz, 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 gmres_device_solve(this, ax, x, f, n, coef, blst, gs_h, niter)
Standard GMRES solve.
subroutine gmres_device_free(this)
Deallocate 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, blstx, blsty, blstz, gs_h, niter)
Standard GMRES coupled solve.
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.
Base type for a matrix-vector product providing .
A list of allocatable `bc_t`. Follows the standard interface of lists.
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.