Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
fusedcg_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
50 use utils, only : neko_error
52 use mpi_f08, only : mpi_allreduce, mpi_in_place, mpi_sum
53 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, &
54 c_associated, c_size_t, c_sizeof, c_int, c_loc
55 implicit none
56 private
57
58 integer, parameter :: device_fusedcg_p_space = 10
59
61 type, public, extends(ksp_t) :: fusedcg_device_t
62 real(kind=rp), allocatable :: w(:)
63 real(kind=rp), allocatable :: r(:)
64 real(kind=rp), allocatable :: z(:)
65 real(kind=rp), allocatable :: p(:,:)
66 real(kind=rp), allocatable :: alpha(:)
67 type(c_ptr) :: w_d = c_null_ptr
68 type(c_ptr) :: r_d = c_null_ptr
69 type(c_ptr) :: z_d = c_null_ptr
70 type(c_ptr) :: alpha_d = c_null_ptr
71 type(c_ptr) :: p_d_d = c_null_ptr
72 type(c_ptr), allocatable :: p_d(:)
73 type(c_ptr) :: gs_event = c_null_ptr
74 contains
75 procedure, pass(this) :: init => fusedcg_device_init
76 procedure, pass(this) :: free => fusedcg_device_free
77 procedure, pass(this) :: solve => fusedcg_device_solve
78 procedure, pass(this) :: solve_coupled => fusedcg_device_solve_coupled
79 end type fusedcg_device_t
80
81#ifdef HAVE_CUDA
82 interface
83 subroutine cuda_fusedcg_update_p(p_d, z_d, po_d, beta, n) &
84 bind(c, name = 'cuda_fusedcg_update_p')
85 use, intrinsic :: iso_c_binding
86 import c_rp
87 implicit none
88 type(c_ptr), value :: p_d, z_d, po_d
89 real(c_rp) :: beta
90 integer(c_int) :: n
91 end subroutine cuda_fusedcg_update_p
92 end interface
93
94 interface
95 subroutine cuda_fusedcg_update_x(x_d, p_d, alpha, p_cur, n) &
96 bind(c, name = 'cuda_fusedcg_update_x')
97 use, intrinsic :: iso_c_binding
98 implicit none
99 type(c_ptr), value :: x_d, p_d, alpha
100 integer(c_int) :: p_cur, n
101 end subroutine cuda_fusedcg_update_x
102 end interface
103
104 interface
105 real(c_rp) function cuda_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, &
106 p_cur, n) bind(c, name = 'cuda_fusedcg_part2')
107 use, intrinsic :: iso_c_binding
108 import c_rp
109 implicit none
110 type(c_ptr), value :: a_d, b_d, c_d, alpha_d
111 real(c_rp) :: alpha
112 integer(c_int) :: n, p_cur
113 end function cuda_fusedcg_part2
114 end interface
115#elif HAVE_HIP
116 interface
117 subroutine hip_fusedcg_update_p(p_d, z_d, po_d, beta, n) &
118 bind(c, name = 'hip_fusedcg_update_p')
119 use, intrinsic :: iso_c_binding
120 import c_rp
121 implicit none
122 type(c_ptr), value :: p_d, z_d, po_d
123 real(c_rp) :: beta
124 integer(c_int) :: n
125 end subroutine hip_fusedcg_update_p
126 end interface
127
128 interface
129 subroutine hip_fusedcg_update_x(x_d, p_d, alpha, p_cur, n) &
130 bind(c, name = 'hip_fusedcg_update_x')
131 use, intrinsic :: iso_c_binding
132 implicit none
133 type(c_ptr), value :: x_d, p_d, alpha
134 integer(c_int) :: p_cur, n
135 end subroutine hip_fusedcg_update_x
136 end interface
137
138 interface
139 real(c_rp) function hip_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, &
140 p_cur, n) bind(c, name = 'hip_fusedcg_part2')
141 use, intrinsic :: iso_c_binding
142 import c_rp
143 implicit none
144 type(c_ptr), value :: a_d, b_d, c_d, alpha_d
145 real(c_rp) :: alpha
146 integer(c_int) :: n, p_cur
147 end function hip_fusedcg_part2
148 end interface
149#endif
150
151contains
152
153 subroutine device_fusedcg_update_p(p_d, z_d, po_d, beta, n)
154 type(c_ptr), value :: p_d, z_d, po_d
155 real(c_rp) :: beta
156 integer(c_int) :: n
157#ifdef HAVE_HIP
158 call hip_fusedcg_update_p(p_d, z_d, po_d, beta, n)
159#elif HAVE_CUDA
160 call cuda_fusedcg_update_p(p_d, z_d, po_d, beta, n)
161#else
162 call neko_error('No device backend configured')
163#endif
164 end subroutine device_fusedcg_update_p
165
166 subroutine device_fusedcg_update_x(x_d, p_d, alpha, p_cur, n)
167 type(c_ptr), value :: x_d, p_d, alpha
168 integer(c_int) :: p_cur, n
169#ifdef HAVE_HIP
170 call hip_fusedcg_update_x(x_d, p_d, alpha, p_cur, n)
171#elif HAVE_CUDA
172 call cuda_fusedcg_update_x(x_d, p_d, alpha, p_cur, n)
173#else
174 call neko_error('No device backend configured')
175#endif
176 end subroutine device_fusedcg_update_x
177
178 function device_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, &
179 p_cur, n) result(res)
180 type(c_ptr), value :: a_d, b_d, c_d, alpha_d
181 real(c_rp) :: alpha
182 integer :: n, p_cur
183 real(kind=rp) :: res
184 integer :: ierr
185#ifdef HAVE_HIP
186 res = hip_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, p_cur, n)
187#elif HAVE_CUDA
188 res = cuda_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, p_cur, n)
189#else
190 call neko_error('No device backend configured')
191#endif
192
193#ifndef HAVE_DEVICE_MPI
194 if (pe_size .gt. 1) then
195 call mpi_allreduce(mpi_in_place, res, 1, &
196 mpi_real_precision, mpi_sum, neko_comm, ierr)
197 end if
198#endif
199
200 end function device_fusedcg_part2
201
203 subroutine fusedcg_device_init(this, n, max_iter, M, rel_tol, abs_tol, &
204 monitor)
205 class(fusedcg_device_t), target, intent(inout) :: this
206 class(pc_t), optional, intent(in), target :: M
207 integer, intent(in) :: n
208 integer, intent(in) :: max_iter
209 real(kind=rp), optional, intent(in) :: rel_tol
210 real(kind=rp), optional, intent(in) :: abs_tol
211 logical, optional, intent(in) :: monitor
212 type(c_ptr) :: ptr
213 integer(c_size_t) :: p_size
214 integer :: i
215
216 call this%free()
217
218 allocate(this%w(n))
219 allocate(this%r(n))
220 allocate(this%z(n))
221 allocate(this%p(n, device_fusedcg_p_space))
222 allocate(this%p_d(device_fusedcg_p_space))
223 allocate(this%alpha(device_fusedcg_p_space))
224
225 if (present(m)) then
226 this%M => m
227 end if
228
229 call device_map(this%w, this%w_d, n)
230 call device_map(this%r, this%r_d, n)
231 call device_map(this%z, this%z_d, n)
232 call device_map(this%alpha, this%alpha_d, device_fusedcg_p_space)
233 do i = 1, device_fusedcg_p_space
234 this%p_d(i) = c_null_ptr
235 call device_map(this%p(:,i), this%p_d(i), n)
236 end do
237
238 p_size = c_sizeof(c_null_ptr) * (device_fusedcg_p_space)
239 call device_alloc(this%p_d_d, p_size)
240 ptr = c_loc(this%p_d)
241 call device_memcpy(ptr, this%p_d_d, p_size, &
242 host_to_device, sync = .false.)
243 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
244 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
245 else if (present(rel_tol) .and. present(abs_tol)) then
246 call this%ksp_init(max_iter, rel_tol, abs_tol)
247 else if (present(monitor) .and. present(abs_tol)) then
248 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
249 else if (present(rel_tol) .and. present(monitor)) then
250 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
251 else if (present(rel_tol)) then
252 call this%ksp_init(max_iter, rel_tol = rel_tol)
253 else if (present(abs_tol)) then
254 call this%ksp_init(max_iter, abs_tol = abs_tol)
255 else if (present(monitor)) then
256 call this%ksp_init(max_iter, monitor = monitor)
257 else
258 call this%ksp_init(max_iter)
259 end if
260
261 call device_event_create(this%gs_event, 2)
262
263 end subroutine fusedcg_device_init
264
266 subroutine fusedcg_device_free(this)
267 class(fusedcg_device_t), intent(inout) :: this
268 integer :: i
269
270 call this%ksp_free()
271
272 if (allocated(this%w)) then
273 if (c_associated(this%w_d)) then
274 call device_unmap(this%w, this%w_d)
275 end if
276 deallocate(this%w)
277 end if
278
279 if (allocated(this%r)) then
280 if (c_associated(this%r_d)) then
281 call device_unmap(this%r, this%r_d)
282 end if
283 deallocate(this%r)
284 end if
285
286
287 if (allocated(this%z)) then
288 if (c_associated(this%z_d)) then
289 call device_unmap(this%z, this%z_d)
290 end if
291 deallocate(this%z)
292 end if
293
294
295 if (allocated(this%alpha)) then
296 if (c_associated(this%alpha_d)) then
297 call device_unmap(this%alpha, this%alpha_d)
298 end if
299 deallocate(this%alpha)
300 end if
301
302 if (allocated(this%p)) then
303 if (allocated(this%p_d)) then
304 do i = 1, device_fusedcg_p_space
305 if (c_associated(this%p_d(i))) then
306 call device_unmap(this%p(:,i), this%p_d(i))
307 end if
308 end do
309 end if
310 deallocate(this%p)
311 end if
312
313 if (allocated(this%p_d)) then
314 deallocate(this%p_d)
315 end if
316
317 if (c_associated(this%p_d_d)) then
318 call device_free(this%p_d_d)
319 end if
320
321 nullify(this%M)
322
323 if (c_associated(this%gs_event)) then
324 call device_event_destroy(this%gs_event)
325 end if
326
327 end subroutine fusedcg_device_free
328
330 function fusedcg_device_solve(this, Ax, x, f, n, coef, bc_projector, gs_h, &
331 niter) result(ksp_results)
332 class(fusedcg_device_t), intent(inout) :: this
333 class(ax_t), intent(in) :: ax
334 type(field_t), intent(inout) :: x
335 integer, intent(in) :: n
336 real(kind=rp), dimension(n), intent(in) :: f
337 type(coef_t), intent(inout) :: coef
338 class(scalar_bc_projector_t), intent(inout) :: bc_projector
339 type(gs_t), intent(inout) :: gs_h
340 type(ksp_monitor_t) :: ksp_results
341 integer, optional, intent(in) :: niter
342 integer :: iter, max_iter, ierr, i, p_cur, p_prev
343 real(kind=rp) :: rnorm, rtr, norm_fac, rtz1, rtz2
344 real(kind=rp) :: pap, beta
345 type(c_ptr) :: f_d
346 f_d = device_get_ptr(f)
347
348 if (present(niter)) then
349 max_iter = niter
350 else
351 max_iter = ksp_max_iter
352 end if
353 norm_fac = 1.0_rp / sqrt(coef%volume)
354
355 associate(w => this%w, r => this%r, p => this%p, z => this%z, &
356 alpha => this%alpha, alpha_d => this%alpha_d, &
357 w_d => this%w_d, r_d => this%r_d, z_d => this%z_d, &
358 p_d => this%p_d, p_d_d => this%p_d_d)
359
360 rtz1 = 1.0_rp
362 p_cur = 1
363 call device_rzero(x%x_d, n)
364 call device_rzero(p_d(1), n)
365 call device_copy(r_d, f_d, n)
366
367 rtr = device_glsc3(r_d, coef%mult_d, r_d, n)
368 rnorm = sqrt(rtr)*norm_fac
369 ksp_results%res_start = rnorm
370 ksp_results%res_final = rnorm
371 ksp_results%iter = 0
372 if (abscmp(rnorm, 0.0_rp)) then
373 ksp_results%converged = .true.
374 return
375 end if
376
377 call this%monitor_start('FusedCG')
378 do iter = 1, max_iter
379 call this%M%solve(z, r, n)
380 rtz2 = rtz1
381 rtz1 = device_glsc3(r_d, coef%mult_d, z_d, n)
382 beta = rtz1 / rtz2
383 if (iter .eq. 1) beta = 0.0_rp
384 call device_fusedcg_update_p(p_d(p_cur), z_d, p_d(p_prev), beta, n)
385
386 call ax%compute(w, p(1, p_cur), coef, x%msh, x%Xh)
387 call gs_h%op(w, n, gs_op_add, this%gs_event)
388 call device_event_sync(this%gs_event)
389 call bc_projector%apply(w, n)
390
391 pap = device_glsc3(w_d, coef%mult_d, this%p_d(p_cur), n)
392
393 alpha(p_cur) = rtz1 / pap
394 rtr = device_fusedcg_part2(r_d, coef%mult_d, w_d, &
395 alpha_d, alpha(p_cur), p_cur, n)
396 rnorm = sqrt(rtr)*norm_fac
397 call this%monitor_iter(iter, rnorm)
398 if ((p_cur .eq. device_fusedcg_p_space) .or. &
399 (rnorm .lt. this%abs_tol) .or. iter .eq. max_iter) then
400 call device_fusedcg_update_x(x%x_d, p_d_d, alpha_d, p_cur, n)
401 p_prev = p_cur
402 p_cur = 1
403 if (rnorm .lt. this%abs_tol) exit
404 else
405 p_prev = p_cur
406 p_cur = p_cur + 1
407 end if
408 end do
409 call this%monitor_stop()
410 ksp_results%res_final = rnorm
411 ksp_results%iter = iter
412 ksp_results%converged = this%is_converged(iter, rnorm)
413
414 end associate
415
416 end function fusedcg_device_solve
417
419 function fusedcg_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
420 n, coef, bc_projector, gs_h, niter) result(ksp_results)
421 class(fusedcg_device_t), intent(inout) :: this
422 class(ax_t), intent(in) :: ax
423 type(field_t), intent(inout) :: x
424 type(field_t), intent(inout) :: y
425 type(field_t), intent(inout) :: z
426 integer, intent(in) :: n
427 real(kind=rp), dimension(n), intent(in) :: fx
428 real(kind=rp), dimension(n), intent(in) :: fy
429 real(kind=rp), dimension(n), intent(in) :: fz
430 type(coef_t), intent(inout) :: coef
431 class(vector_bc_projector_t), intent(inout) :: bc_projector
432 type(gs_t), intent(inout) :: gs_h
433 type(ksp_monitor_t), dimension(3) :: ksp_results
434 integer, optional, intent(in) :: niter
435 type(scalar_bc_projector_t), pointer :: bc_x, bc_y, bc_z
436
437 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
438 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
439 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
440 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
441
443
444end module fusedcg_device
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
void hip_fusedcg_update_x(void *x, void *p, void *alpha, int *p_cur, int *n)
real hip_fusedcg_part2(void *a, void *b, void *c, void *alpha_d, real *alpha, int *p_cur, int *n)
void hip_fusedcg_update_p(void *p, void *z, void *po, real *beta, int *n)
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_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
Defines a fused Conjugate Gradient method for accelerators.
subroutine fusedcg_device_free(this)
Deallocate a pipelined PCG solver.
real(kind=rp) function device_fusedcg_part2(a_d, b_d, c_d, alpha_d, alpha, p_cur, n)
integer, parameter device_fusedcg_p_space
subroutine device_fusedcg_update_x(x_d, p_d, alpha, p_cur, n)
type(ksp_monitor_t) function fusedcg_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Pipelined PCG solve.
subroutine fusedcg_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a fused PCG solver.
subroutine device_fusedcg_update_p(p_d, z_d, po_d, beta, n)
type(ksp_monitor_t) function, dimension(3) fusedcg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Pipelined PCG solve coupled solve.
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
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.
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
Fused preconditioned conjugate gradient method.
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
Defines a canonical Krylov preconditioner.
Definition precon.f90:40
Projector for scalar boundary conditions.
Abstract type for resolving vector boundary conditions.