57 use json_module,
only : json_file
61 use mpi_f08,
only : mpi_in_place, mpi_sum, mpi_allreduce
62 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_loc, &
63 c_associated, c_int, c_size_t, c_sizeof
70 real(kind=
rp),
allocatable :: w(:)
71 real(kind=
rp),
allocatable :: c(:)
72 real(kind=
rp),
allocatable :: r(:)
73 real(kind=
rp),
allocatable :: z(:,:)
74 real(kind=
rp),
allocatable :: h(:,:)
75 real(kind=
rp),
allocatable :: v(:,:)
76 real(kind=
rp),
allocatable :: s(:)
77 real(kind=
rp),
allocatable :: gam(:)
78 type(c_ptr) :: w_d = c_null_ptr
79 type(c_ptr) :: c_d = c_null_ptr
80 type(c_ptr) :: r_d = c_null_ptr
81 type(c_ptr) :: s_d = c_null_ptr
82 type(c_ptr) :: gam_d = c_null_ptr
83 type(c_ptr),
allocatable :: z_d(:), h_d(:), v_d(:)
84 type(c_ptr) :: z_d_d = c_null_ptr
85 type(c_ptr) :: h_d_d = c_null_ptr
86 type(c_ptr) :: v_d_d = c_null_ptr
87 type(c_ptr) :: gs_event = c_null_ptr
99 bind(c, name =
'hip_gmres_part2')
100 use,
intrinsic :: iso_c_binding
103 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
104 integer(c_int) :: j, n
111 bind(c, name =
'cuda_gmres_part2')
112 use,
intrinsic :: iso_c_binding
115 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
116 integer(c_int) :: j, n
124 type(c_ptr),
value :: h_d, w_d, v_d_d, mult_d
125 integer(c_int) :: j, n
133 call neko_error(
'No device backend configured')
136#ifndef HAVE_DEVICE_MPI
137 if (pe_size .gt. 1)
then
138 call mpi_allreduce(mpi_in_place, alpha, 1, &
139 mpi_real_precision, mpi_sum, neko_comm, ierr)
148 integer,
intent(in) :: n
149 integer,
intent(in) :: max_iter
150 class(pc_t),
optional,
intent(in),
target :: M
151 real(kind=rp),
optional,
intent(in) :: rel_tol
152 real(kind=rp),
optional,
intent(in) :: abs_tol
153 logical,
optional,
intent(in) :: monitor
154 type(device_ident_t),
target :: M_ident
156 integer(c_size_t) :: z_size
169 call device_map(this%w, this%w_d, n)
170 call device_map(this%r, this%r_d, n)
172 allocate(this%c(this%m_restart))
173 allocate(this%s(this%m_restart))
174 allocate(this%gam(this%m_restart + 1))
175 call device_map(this%c, this%c_d, this%m_restart)
176 call device_map(this%s, this%s_d, this%m_restart)
177 call device_map(this%gam, this%gam_d, this%m_restart+1)
179 allocate(this%z(n, this%m_restart))
180 allocate(this%v(n, this%m_restart))
181 allocate(this%h(this%m_restart, this%m_restart))
182 allocate(this%z_d(this%m_restart))
183 allocate(this%v_d(this%m_restart))
184 allocate(this%h_d(this%m_restart))
185 do i = 1, this%m_restart
186 this%z_d(i) = c_null_ptr
187 call device_map(this%z(:,i), this%z_d(i), n)
189 this%v_d(i) = c_null_ptr
190 call device_map(this%v(:,i), this%v_d(i), n)
192 this%h_d(i) = c_null_ptr
193 call device_map(this%h(:,i), this%h_d(i), this%m_restart)
196 z_size = c_sizeof(c_null_ptr) * (this%m_restart)
197 call device_alloc(this%z_d_d, z_size)
198 call device_alloc(this%v_d_d, z_size)
199 call device_alloc(this%h_d_d, z_size)
200 ptr = c_loc(this%z_d)
201 call device_memcpy(ptr, this%z_d_d, z_size, &
202 host_to_device, sync = .false.)
203 ptr = c_loc(this%v_d)
204 call device_memcpy(ptr, this%v_d_d, z_size, &
205 host_to_device, sync = .false.)
206 ptr = c_loc(this%h_d)
207 call device_memcpy(ptr, this%h_d_d, z_size, &
208 host_to_device, sync = .false.)
211 if (
present(rel_tol) .and.
present(abs_tol) .and.
present(monitor))
then
212 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
213 else if (
present(rel_tol) .and.
present(abs_tol))
then
214 call this%ksp_init(max_iter, rel_tol, abs_tol)
215 else if (
present(monitor) .and.
present(abs_tol))
then
216 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
217 else if (
present(rel_tol) .and.
present(monitor))
then
218 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
219 else if (
present(rel_tol))
then
220 call this%ksp_init(max_iter, rel_tol = rel_tol)
221 else if (
present(abs_tol))
then
222 call this%ksp_init(max_iter, abs_tol = abs_tol)
223 else if (
present(monitor))
then
224 call this%ksp_init(max_iter, monitor = monitor)
226 call this%ksp_init(max_iter)
229 call device_event_create(this%gs_event, 2)
238 integer,
intent(in) :: n
239 type(json_file),
intent(inout) :: json
240 class(pc_t),
optional,
intent(in),
target :: M
241 character(len=LOG_SIZE) :: log_buf
243 call json_get_or_lookup_or_default(json,
'space_size', &
244 this%m_restart, gmres_space_size)
245 if (this%m_restart .lt. 1)
then
246 call neko_error(
'GMRES space_size must be at least 1')
248 write(log_buf,
'(A,I0)')
'GMRES space: ', this%m_restart
249 call neko_log%message(log_buf)
251 call this%ksp_init_from_json(n, json, m = m)
262 if (
allocated(this%w))
then
263 if (c_associated(this%w_d))
then
264 call device_unmap(this%w, this%w_d)
269 if (
allocated(this%c))
then
270 if (c_associated(this%c_d))
then
271 call device_unmap(this%c, this%c_d)
276 if (
allocated(this%r))
then
277 if (c_associated(this%r_d))
then
278 call device_unmap(this%r, this%r_d)
283 if (
allocated(this%z))
then
284 if (
allocated(this%z_d))
then
285 do i = 1, this%m_restart
286 if (c_associated(this%z_d(i)))
then
287 call device_unmap(this%z(:,i), this%z_d(i))
294 if (
allocated(this%h))
then
295 if (
allocated(this%h_d))
then
296 do i = 1, this%m_restart
297 if (c_associated(this%h_d(i)))
then
298 call device_unmap(this%h(:,i), this%h_d(i))
305 if (
allocated(this%v))
then
306 if (
allocated(this%v_d))
then
307 do i = 1, this%m_restart
308 if (c_associated(this%v_d(i)))
then
309 call device_unmap(this%v(:,i), this%v_d(i))
316 if (
allocated(this%s))
then
317 if (c_associated(this%s_d))
then
318 call device_unmap(this%s, this%s_d)
322 if (
allocated(this%gam))
then
323 if (c_associated(this%gam_d))
then
324 call device_unmap(this%gam, this%gam_d)
329 if (c_associated(this%z_d_d))
then
330 call device_free(this%z_d_d)
332 if (c_associated(this%v_d_d))
then
333 call device_free(this%v_d_d)
335 if (c_associated(this%h_d_d))
then
336 call device_free(this%h_d_d)
341 if (c_associated(this%gs_event))
then
342 call device_event_destroy(this%gs_event)
349 niter)
result(ksp_results)
351 class(ax_t),
intent(in) :: ax
352 type(field_t),
intent(inout) :: x
353 integer,
intent(in) :: n
354 real(kind=rp),
dimension(n),
intent(in) :: f
355 type(coef_t),
intent(inout) :: coef
356 class(scalar_bc_projector_t),
intent(inout) :: bc_projector
357 type(gs_t),
intent(inout) :: gs_h
358 type(ksp_monitor_t) :: ksp_results
359 integer,
optional,
intent(in) :: niter
360 integer :: iter, max_iter
362 real(kind=rp) :: rnorm, alpha, temp, lr, alpha2, norm_fac
366 f_d = device_get_ptr(f)
372 if (
present(niter))
then
375 max_iter = this%max_iter
378 associate(w => this%w, c => this%c, r => this%r, z => this%z, h => this%h, &
379 v => this%v, s => this%s, gam => this%gam, v_d => this%v_d, &
380 w_d => this%w_d, r_d => this%r_d, h_d => this%h_d, &
381 v_d_d => this%v_d_d, x_d => x%x_d, z_d_d => this%z_d_d, &
384 norm_fac = 1.0_rp / sqrt(coef%volume)
385 call rzero(gam, this%m_restart + 1)
386 call rone(s, this%m_restart)
387 call rone(c, this%m_restart)
388 call rzero(h, this%m_restart * this%m_restart)
389 call device_rzero(x%x_d, n)
390 call device_rzero(this%gam_d, this%m_restart + 1)
391 call device_rone(this%s_d, this%m_restart)
392 call device_rone(this%c_d, this%m_restart)
394 call rzero(this%h, this%m_restart**2)
399 call this%monitor_start(
'GMRES')
400 do while (.not. conv .and. iter .lt. max_iter)
402 if (iter .eq. 0)
then
403 call device_copy(r_d, f_d, n)
405 call device_copy(r_d, f_d, n)
406 call ax%compute(w, x%x, coef, x%msh, x%Xh)
407 call gs_h%op(w, n, gs_op_add, this%gs_event)
408 call device_event_sync(this%gs_event)
409 call bc_projector%apply(w, n)
410 call device_sub2(r_d, w_d, n)
413 gam(1) = sqrt(device_glsc3(r_d, r_d, coef%mult_d, n))
414 if (iter .eq. 0)
then
415 ksp_results%res_start = gam(1) * norm_fac
418 if (abscmp(gam(1), 0.0_rp))
exit
421 temp = 1.0_rp / gam(1)
422 call device_cmult2(v_d(1), r_d, temp, n)
423 do j = 1, this%m_restart
426 call this%M%solve(z(1,j), v(1,j), n)
428 call ax%compute(w, z(1,j), coef, x%msh, x%Xh)
429 call gs_h%op(w, n, gs_op_add, this%gs_event)
430 call device_event_sync(this%gs_event)
431 call bc_projector%apply(w, n)
433 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
435 h(i,j) = device_glsc3(w_d, v_d(i), coef%mult_d, n)
437 call device_add2s2(w_d, v_d(i), -h(i,j), n)
439 alpha2 = device_glsc3(w_d, w_d, coef%mult_d, n)
442 call device_glsc3_many(h(1,j), w_d, v_d_d, coef%mult_d, j, n)
444 call device_memcpy(h(:,j), h_d(j), j, &
445 host_to_device, sync = .false.)
455 h(i,j) = c(i)*temp + s(i) * h(i+1,j)
456 h(i+1,j) = -s(i)*temp + c(i) * h(i+1,j)
460 if (abscmp(alpha, 0.0_rp))
then
465 lr = sqrt(h(j,j) * h(j,j) + alpha2)
470 call device_memcpy(h(:,j), h_d(j), j, &
471 host_to_device, sync = .false.)
472 gam(j+1) = -s(j) * gam(j)
473 gam(j) = c(j) * gam(j)
475 rnorm = abs(gam(j+1)) * norm_fac
476 call this%monitor_iter(iter, rnorm)
477 if (rnorm .lt. this%abs_tol)
then
482 if (iter + 1 .gt. max_iter)
exit
484 if (j .lt. this%m_restart)
then
485 temp = 1.0_rp / alpha
486 call device_cmult2(v_d(j+1), w_d, temp, n)
491 j = min(j, this%m_restart)
495 temp = temp - h(k,i) * c(i)
500 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1)
then
502 call device_add2s2(x_d, this%z_d(i), c(i), n)
505 call device_memcpy(c, c_d, j, host_to_device, sync = .false.)
506 call device_add2s2_many(x_d, z_d_d, c_d, j, n)
511 call this%monitor_stop()
512 ksp_results%res_final = rnorm
513 ksp_results%iter = iter
514 ksp_results%converged = this%is_converged(iter, rnorm)
520 n, coef, bc_projector, gs_h, niter)
result(ksp_results)
522 class(ax_t),
intent(in) :: ax
523 type(field_t),
intent(inout) :: x
524 type(field_t),
intent(inout) :: y
525 type(field_t),
intent(inout) :: z
526 integer,
intent(in) :: n
527 real(kind=rp),
dimension(n),
intent(in) :: fx
528 real(kind=rp),
dimension(n),
intent(in) :: fy
529 real(kind=rp),
dimension(n),
intent(in) :: fz
530 type(coef_t),
intent(inout) :: coef
531 class(vector_bc_projector_t),
intent(inout) :: bc_projector
532 type(gs_t),
intent(inout) :: gs_h
533 type(ksp_monitor_t),
dimension(3) :: ksp_results
534 integer,
optional,
intent(in) :: niter
535 type(scalar_bc_projector_t),
pointer :: bc_x, bc_y, bc_z
537 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
538 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
539 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
540 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.
subroutine gmres_device_init_from_json(this, n, json, m)
Initialise a standard GMRES solver from a JSON solver block. Reads space_size (default GMRES_SPACE_SI...
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.
integer, parameter, public gmres_space_size
Default Krylov space size (restart length) of the GMRES solvers.
Utilities for retrieving parameters from the case files.
Implements the base abstract type for Krylov solvers plus helper types.
type(log_t), public neko_log
Global log stream.
integer, parameter, public log_size
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.