Neko 1.99.6
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
gmres_device.F90
Go to the documentation of this file.
1! Copyright (c) 2022-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 krylov, only : ksp_t, ksp_monitor_t
37 use precon, only : pc_t
38 use ax_product, only : ax_t
39 use num_types, only : rp, c_rp
40 use field, only : field_t
41 use coefs, only : coef_t
42 use gather_scatter, only : gs_t, gs_op_add
43 use bc_list, only : bc_list_t
45 use math, only : rone, rzero, abscmp
53 use utils, only : neko_error
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
58 implicit none
59 private
60
62 type, public, extends(ksp_t) :: gmres_device_t
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
82 contains
83 procedure, pass(this) :: init => gmres_device_init
84 procedure, pass(this) :: free => gmres_device_free
85 procedure, pass(this) :: solve => gmres_device_solve
86 procedure, pass(this) :: solve_coupled => gmres_device_solve_coupled
87 end type gmres_device_t
88
89#ifdef HAVE_HIP
90 interface
91 real(c_rp) function hip_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) &
92 bind(c, name = 'hip_gmres_part2')
93 use, intrinsic :: iso_c_binding
94 import c_rp
95 implicit none
96 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
97 integer(c_int) :: j, n
98 end function hip_gmres_part2
99 end interface
100#elif HAVE_CUDA
101
102 interface
103 real(c_rp) function cuda_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) &
104 bind(c, name = 'cuda_gmres_part2')
105 use, intrinsic :: iso_c_binding
106 import c_rp
107 implicit none
108 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
109 integer(c_int) :: j, n
110 end function cuda_gmres_part2
111 end interface
112#endif
113
114contains
115
116 function device_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n) result(alpha)
117 type(c_ptr), value :: h_d, w_d, v_d_d, mult_d
118 integer(c_int) :: j, n
119 real(c_rp) :: alpha
120 integer :: ierr
121#ifdef HAVE_HIP
122 alpha = hip_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
123#elif HAVE_CUDA
124 alpha = cuda_gmres_part2(w_d, v_d_d, h_d, mult_d, j, n)
125#else
126 call neko_error('No device backend configured')
127#endif
128
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)
133 end if
134#endif
135
136 end function device_gmres_part2
137
139 subroutine gmres_device_init(this, n, max_iter, M, rel_tol, abs_tol, monitor)
140 class(gmres_device_t), target, intent(inout) :: this
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
148 type(c_ptr) :: ptr
149 integer(c_size_t) :: z_size
150 integer :: i
151
152 call this%free()
153
154 if (present(m)) then
155 this%M => m
156 else
157 this%M => m_ident
158 end if
159
160 allocate(this%w(n))
161 allocate(this%r(n))
162 call device_map(this%w, this%w_d, n)
163 call device_map(this%r, this%r_d, n)
164
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)
171
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)
181
182 this%v_d(i) = c_null_ptr
183 call device_map(this%v(:,i), this%v_d(i), n)
184
185 this%h_d(i) = c_null_ptr
186 call device_map(this%h(:,i), this%h_d(i), this%m_restart)
187 end do
188
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.)
202
203
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)
218 else
219 call this%ksp_init(max_iter)
220 end if
221
222 call device_event_create(this%gs_event, 2)
223
224 end subroutine gmres_device_init
225
227 subroutine gmres_device_free(this)
228 class(gmres_device_t), intent(inout) :: this
229 integer :: i
230
231 call this%ksp_free()
232
233 if (allocated(this%w)) then
234 if (c_associated(this%w_d)) then
235 call device_unmap(this%w, this%w_d)
236 end if
237 deallocate(this%w)
238 end if
239
240 if (allocated(this%c)) then
241 if (c_associated(this%c_d)) then
242 call device_unmap(this%c, this%c_d)
243 end if
244 deallocate(this%c)
245 end if
246
247 if (allocated(this%r)) then
248 if (c_associated(this%r_d)) then
249 call device_unmap(this%r, this%r_d)
250 end if
251 deallocate(this%r)
252 end if
253
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))
259 end if
260 end do
261 end if
262 deallocate(this%z)
263 end if
264
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))
270 end if
271 end do
272 end if
273 deallocate(this%h)
274 end if
275
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))
281 end if
282 end do
283 end if
284 deallocate(this%v)
285 end if
286
287 if (allocated(this%s)) then
288 if (c_associated(this%s_d)) then
289 call device_unmap(this%s, this%s_d)
290 end if
291 deallocate(this%s)
292 end if
293 if (allocated(this%gam)) then
294 if (c_associated(this%gam_d)) then
295 call device_unmap(this%gam, this%gam_d)
296 end if
297 deallocate(this%gam)
298 end if
299
300 if (c_associated(this%z_d_d)) then
301 call device_free(this%z_d_d)
302 end if
303 if (c_associated(this%v_d_d)) then
304 call device_free(this%v_d_d)
305 end if
306 if (c_associated(this%h_d_d)) then
307 call device_free(this%h_d_d)
308 end if
309
310 nullify(this%M)
311
312 if (c_associated(this%gs_event)) then
313 call device_event_destroy(this%gs_event)
314 end if
315
316 end subroutine gmres_device_free
317
319 function gmres_device_solve(this, Ax, x, f, n, coef, blst, gs_h, niter) &
320 result(ksp_results)
321 class(gmres_device_t), intent(inout) :: this
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
332 integer :: i, j, k
333 real(kind=rp) :: rnorm, alpha, temp, lr, alpha2, norm_fac
334 logical :: conv
335 type(c_ptr) :: f_d
336
337 f_d = device_get_ptr(f)
338
339 conv = .false.
340 iter = 0
341 rnorm = 0.0_rp
342
343 if (present(niter)) then
344 max_iter = niter
345 else
346 max_iter = this%max_iter
347 end if
348
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, &
353 c_d => this%c_d)
354
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)
364
365 call rzero(this%h, this%m_restart**2)
366 ! do j = 1, this%m_restart
367 ! call device_rzero(h_d(j), this%m_restart)
368 ! end do
369
370 call this%monitor_start('GMRES')
371 do while (.not. conv .and. iter .lt. max_iter)
372
373 if (iter .eq. 0) then
374 call device_copy(r_d, f_d, n)
375 else
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)
382 end if
383
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
387 end if
388
389 if (abscmp(gam(1), 0.0_rp)) exit
390
391 rnorm = 0.0_rp
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
395 iter = iter+1
396
397 call this%M%solve(z(1,j), v(1,j), n)
398
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)
403
404 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1) then
405 do i = 1, j
406 h(i,j) = device_glsc3(w_d, v_d(i), coef%mult_d, n)
407
408 call device_add2s2(w_d, v_d(i), -h(i,j), n)
409
410 alpha2 = device_glsc3(w_d, w_d, coef%mult_d, n)
411 end do
412 else
413 call device_glsc3_many(h(1,j), w_d, v_d_d, coef%mult_d, j, n)
414
415 call device_memcpy(h(:,j), h_d(j), j, &
416 host_to_device, sync = .false.)
417
418 alpha2 = device_gmres_part2(w_d, v_d_d, h_d(j), &
419 coef%mult_d, j, n)
420
421 end if
422
423 alpha = sqrt(alpha2)
424 do i = 1, j-1
425 temp = h(i,j)
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)
428 end do
429
430 rnorm = 0.0_rp
431 if (abscmp(alpha, 0.0_rp)) then
432 conv = .true.
433 exit
434 end if
435
436 lr = sqrt(h(j,j) * h(j,j) + alpha2)
437 temp = 1.0_rp / lr
438 c(j) = h(j,j) * temp
439 s(j) = alpha * temp
440 h(j,j) = lr
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)
445
446 rnorm = abs(gam(j+1)) * norm_fac
447 call this%monitor_iter(iter, rnorm)
448 if (rnorm .lt. this%abs_tol) then
449 conv = .true.
450 exit
451 end if
452
453 if (iter + 1 .gt. max_iter) exit
454
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)
458 end if
459
460 end do
461
462 j = min(j, this%m_restart)
463 do k = j, 1, -1
464 temp = gam(k)
465 do i = j, k+1, -1
466 temp = temp - h(k,i) * c(i)
467 end do
468 c(k) = temp / h(k,k)
469 end do
470
471 if (neko_bcknd_opencl .eq. 1 .or. neko_bcknd_metal .eq. 1) then
472 do i = 1, j
473 call device_add2s2(x_d, this%z_d(i), c(i), n)
474 end do
475 else
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)
478 end if
479 end do
480
481 end associate
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)
486
487 end function gmres_device_solve
488
490 function gmres_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
491 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
492 class(gmres_device_t), intent(inout) :: this
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
508
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)
512
513 end function gmres_device_solve_coupled
514
515end module gmres_device
__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)
Definition gmres_aux.cu:62
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
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.
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.
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.
Definition krylov.f90:34
Definition math.f90:60
subroutine, public rone(a, n)
Set all elements to one.
Definition math.f90:277
subroutine, public rzero(a, n)
Zero a real vector.
Definition math.f90:235
Build configurations.
integer, parameter neko_bcknd_opencl
integer, parameter neko_bcknd_metal
integer, parameter, public c_rp
Definition num_types.f90:13
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Krylov preconditioner.
Definition precon.f90:34
Utilities.
Definition utils.f90:35
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
Defines a canonical Krylov preconditioner for accelerators.
Gather-scatter kernel.
Standard preconditioned generalized minimal residual method.
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
Defines a canonical Krylov preconditioner.
Definition precon.f90:40