Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
cg_cpld_device.f90
Go to the documentation of this file.
1! Copyright (c) 2025-2026, 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!
35 use num_types, only : rp
37 use precon, only : pc_t
38 use ax_product, only : ax_t
39 use field, only : field_t
40 use coefs, only : coef_t
41 use gather_scatter, only : gs_t, gs_op_add
44 use math, only : abscmp
50 use utils, only : neko_error
51 use operators, only : rotate_cyc
52 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
53 implicit none
54 private
55
57 type, public, extends(ksp_t) :: cg_cpld_device_t
58 real(kind=rp), allocatable :: w1(:)
59 real(kind=rp), allocatable :: w2(:)
60 real(kind=rp), allocatable :: w3(:)
61 real(kind=rp), allocatable :: r1(:)
62 real(kind=rp), allocatable :: r2(:)
63 real(kind=rp), allocatable :: r3(:)
64 real(kind=rp), allocatable :: p1(:)
65 real(kind=rp), allocatable :: p2(:)
66 real(kind=rp), allocatable :: p3(:)
67 real(kind=rp), allocatable :: z1(:)
68 real(kind=rp), allocatable :: z2(:)
69 real(kind=rp), allocatable :: z3(:)
70 real(kind=rp), allocatable :: tmp(:)
71
72
73 type(c_ptr) :: w1_d = c_null_ptr
74 type(c_ptr) :: w2_d = c_null_ptr
75 type(c_ptr) :: w3_d = c_null_ptr
76
77 type(c_ptr) :: r1_d = c_null_ptr
78 type(c_ptr) :: r2_d = c_null_ptr
79 type(c_ptr) :: r3_d = c_null_ptr
80
81 type(c_ptr) :: p1_d = c_null_ptr
82 type(c_ptr) :: p2_d = c_null_ptr
83 type(c_ptr) :: p3_d = c_null_ptr
84
85 type(c_ptr) :: z1_d = c_null_ptr
86 type(c_ptr) :: z2_d = c_null_ptr
87 type(c_ptr) :: z3_d = c_null_ptr
88
89 type(c_ptr) :: tmp_d = c_null_ptr
90
91 type(c_ptr) :: gs_event = c_null_ptr
92 contains
93 procedure, pass(this) :: init => cg_cpld_device_init
94 procedure, pass(this) :: free => cg_cpld_device_free
95 procedure, pass(this) :: solve => cg_cpld_device_nop
96 procedure, pass(this) :: solve_coupled => cg_cpld_device_solve
97 end type cg_cpld_device_t
98
99contains
100
102 subroutine cg_cpld_device_init(this, n, max_iter, M, rel_tol, abs_tol, &
103 monitor)
104 class(cg_cpld_device_t), target, intent(inout) :: this
105 class(pc_t), optional, intent(in), target :: M
106 integer, intent(in) :: n
107 integer, intent(in) :: max_iter
108 real(kind=rp), optional, intent(in) :: rel_tol
109 real(kind=rp), optional, intent(in) :: abs_tol
110 logical, optional, intent(in) :: monitor
111
112 call this%free()
113
114 allocate(this%w1(n))
115 allocate(this%w2(n))
116 allocate(this%w3(n))
117 allocate(this%r1(n))
118 allocate(this%r2(n))
119 allocate(this%r3(n))
120 allocate(this%p1(n))
121 allocate(this%p2(n))
122 allocate(this%p3(n))
123 allocate(this%z1(n))
124 allocate(this%z2(n))
125 allocate(this%z3(n))
126 allocate(this%tmp(n))
127
128 call device_map(this%tmp, this%tmp_d, n)
129 call device_map(this%z1, this%z1_d, n)
130 call device_map(this%z2, this%z2_d, n)
131 call device_map(this%z3, this%z3_d, n)
132 call device_map(this%p1, this%p1_d, n)
133 call device_map(this%p2, this%p2_d, n)
134 call device_map(this%p3, this%p3_d, n)
135 call device_map(this%r1, this%r1_d, n)
136 call device_map(this%r2, this%r2_d, n)
137 call device_map(this%r3, this%r3_d, n)
138 call device_map(this%w1, this%w1_d, n)
139 call device_map(this%w2, this%w2_d, n)
140 call device_map(this%w3, this%w3_d, n)
141
142 if (present(m)) then
143 this%M => m
144 end if
145
146 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
147 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
148 else if (present(rel_tol) .and. present(abs_tol)) then
149 call this%ksp_init(max_iter, rel_tol, abs_tol)
150 else if (present(monitor) .and. present(abs_tol)) then
151 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
152 else if (present(rel_tol) .and. present(monitor)) then
153 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
154 else if (present(rel_tol)) then
155 call this%ksp_init(max_iter, rel_tol = rel_tol)
156 else if (present(abs_tol)) then
157 call this%ksp_init(max_iter, abs_tol = abs_tol)
158 else if (present(monitor)) then
159 call this%ksp_init(max_iter, monitor = monitor)
160 else
161 call this%ksp_init(max_iter)
162 end if
163
164 call device_event_create(this%gs_event, 2)
165 end subroutine cg_cpld_device_init
166
168 subroutine cg_cpld_device_free(this)
169 class(cg_cpld_device_t), intent(inout) :: this
170
171 call this%ksp_free()
172
173 if (allocated(this%w1)) then
174 if (c_associated(this%w1_d)) then
175 call device_unmap(this%w1, this%w1_d)
176 end if
177 deallocate(this%w1)
178 end if
179
180 if (allocated(this%w2)) then
181 if (c_associated(this%w2_d)) then
182 call device_unmap(this%w2, this%w2_d)
183 end if
184 deallocate(this%w2)
185 end if
186
187 if (allocated(this%w3)) then
188 if (c_associated(this%w3_d)) then
189 call device_unmap(this%w3, this%w3_d)
190 end if
191 deallocate(this%w3)
192 end if
193
194 if (allocated(this%r1)) then
195 if (c_associated(this%r1_d)) then
196 call device_unmap(this%r1, this%r1_d)
197 end if
198 deallocate(this%r1)
199 end if
200
201 if (allocated(this%r2)) then
202 if (c_associated(this%r2_d)) then
203 call device_unmap(this%r2, this%r2_d)
204 end if
205 deallocate(this%r2)
206 end if
207
208 if (allocated(this%r3)) then
209 if (c_associated(this%r3_d)) then
210 call device_unmap(this%r3, this%r3_d)
211 end if
212 deallocate(this%r3)
213 end if
214
215 if (allocated(this%p1)) then
216 if (c_associated(this%p1_d)) then
217 call device_unmap(this%p1, this%p1_d)
218 end if
219 deallocate(this%p1)
220 end if
221
222 if (allocated(this%p2)) then
223 if (c_associated(this%p2_d)) then
224 call device_unmap(this%p2, this%p2_d)
225 end if
226 deallocate(this%p2)
227 end if
228
229 if (allocated(this%p3)) then
230 if (c_associated(this%p3_d)) then
231 call device_unmap(this%p3, this%p3_d)
232 end if
233 deallocate(this%p3)
234 end if
235
236 if (allocated(this%z1)) then
237 if (c_associated(this%z1_d)) then
238 call device_unmap(this%z1, this%z1_d)
239 end if
240 deallocate(this%z1)
241 end if
242
243 if (allocated(this%z2)) then
244 if (c_associated(this%z2_d)) then
245 call device_unmap(this%z2, this%z2_d)
246 end if
247 deallocate(this%z2)
248 end if
249
250 if (allocated(this%z3)) then
251 if (c_associated(this%z3_d)) then
252 call device_unmap(this%z3, this%z3_d)
253 end if
254 deallocate(this%z3)
255 end if
256
257 if (allocated(this%tmp)) then
258 if (c_associated(this%tmp_d)) then
259 call device_unmap(this%tmp, this%tmp_d)
260 end if
261 deallocate(this%tmp)
262 end if
263
264 nullify(this%M)
265
266 if (c_associated(this%gs_event)) then
267 call device_event_destroy(this%gs_event)
268 end if
269
270 end subroutine cg_cpld_device_free
271
272 function cg_cpld_device_nop(this, Ax, x, f, n, coef, bc_projector, gs_h, &
273 niter) result(ksp_results)
274 class(cg_cpld_device_t), intent(inout) :: this
275 class(ax_t), intent(in) :: ax
276 type(field_t), intent(inout) :: x
277 integer, intent(in) :: n
278 real(kind=rp), dimension(n), intent(in) :: f
279 type(coef_t), intent(inout) :: coef
280 class(scalar_bc_projector_t), intent(inout) :: bc_projector
281 type(gs_t), intent(inout) :: gs_h
282 type(ksp_monitor_t) :: ksp_results
283 integer, optional, intent(in) :: niter
284
285 ! Throw and error
286 call neko_error('The cpldcg solver is only defined for coupled solves')
287
288 ksp_results%res_final = 0.0
289 ksp_results%iter = 0
290 end function cg_cpld_device_nop
291
293 function cg_cpld_device_solve(this, Ax, x, y, z, fx, fy, fz, &
294 n, coef, bc_projector, gs_h, niter) result(ksp_results)
295 class(cg_cpld_device_t), intent(inout) :: this
296 class(ax_t), intent(in) :: ax
297 type(field_t), intent(inout) :: x
298 type(field_t), intent(inout) :: y
299 type(field_t), intent(inout) :: z
300 integer, intent(in) :: n
301 real(kind=rp), dimension(n), intent(in) :: fx
302 real(kind=rp), dimension(n), intent(in) :: fy
303 real(kind=rp), dimension(n), intent(in) :: fz
304 type(coef_t), intent(inout) :: coef
305 class(vector_bc_projector_t), intent(inout) :: bc_projector
306 type(gs_t), intent(inout) :: gs_h
307 type(ksp_monitor_t), dimension(3) :: ksp_results
308 integer, optional, intent(in) :: niter
309 integer :: i, iter, max_iter
310 real(kind=rp) :: rnorm, rtr, rtr0, rtz2, rtz1
311 real(kind=rp) :: beta, pap, alpha, alphm, norm_fac
312 integer, parameter :: gdim = 3
313 type(c_ptr) :: fx_d
314 type(c_ptr) :: fy_d
315 type(c_ptr) :: fz_d
316
317 fx_d = device_get_ptr(fx)
318 fy_d = device_get_ptr(fy)
319 fz_d = device_get_ptr(fz)
320
321 if (present(niter)) then
322 max_iter = niter
323 else
324 max_iter = this%max_iter
325 end if
326 norm_fac = 1.0_rp / sqrt(coef%volume)
327
328 associate(p1_d => this%p1_d, p2_d => this%p2_d, p3_d => this%p3_d, &
329 z1_d => this%z1_d, z2_d => this%z2_d, z3_d => this%z3_d, &
330 r1_d => this%r1_d, r2_d => this%r2_d, r3_d => this%r3_d, &
331 w1_d => this%w1_d, w2_d => this%w2_d, w3_d => this%w3_d, &
332 tmp_d => this%tmp_d)
333
334 rtz1 = 1.0_rp
335 call device_rzero(x%x_d, n)
336 call device_rzero(y%x_d, n)
337 call device_rzero(z%x_d, n)
338 call device_rzero(p1_d, n)
339 call device_rzero(p2_d, n)
340 call device_rzero(p3_d, n)
341 call device_rzero(z1_d, n)
342 call device_rzero(z2_d, n)
343 call device_rzero(z3_d, n)
344 call device_copy(r1_d, fx_d, n)
345 call device_copy(r2_d, fy_d, n)
346 call device_copy(r3_d, fz_d, n)
347 call device_vdot3(tmp_d, r1_d, r2_d, r3_d, r1_d, r2_d, r3_d, n)
348
349
350 rtr = device_glsc2(tmp_d, coef%mult_d, n)
351 rnorm = sqrt(rtr)*norm_fac
352 ksp_results%res_start = rnorm
353 ksp_results%res_final = rnorm
354 ksp_results%iter = 0
355 if (abscmp(rnorm, 0.0_rp)) then
356 ksp_results%converged = .true.
357 return
358 end if
359
360 call this%monitor_start('device_cpldCG')
361 do iter = 1, max_iter
362 call this%M%solve(this%z1, this%r1, n)
363 call this%M%solve(this%z2, this%r2, n)
364 call this%M%solve(this%z3, this%r3, n)
365 rtz2 = rtz1
366
367 call device_vdot3(tmp_d, z1_d, z2_d, z3_d, r1_d, r2_d, r3_d, n)
368
369 rtz1 = device_glsc2(tmp_d, coef%mult_d, n)
370
371 beta = rtz1 / rtz2
372 if (iter .eq. 1) beta = 0.0_rp
373 call device_add2s1(p1_d, z1_d, beta, n)
374 call device_add2s1(p2_d, z2_d, beta, n)
375 call device_add2s1(p3_d, z3_d, beta, n)
376
377 call ax%compute_vector(this%w1, this%w2, this%w3, &
378 this%p1, this%p2, this%p3, coef, x%msh, x%Xh)
379
380 call rotate_cyc(w1_d, w2_d, w3_d, 1, coef)
381 call gs_h%op(this%w1, this%w2, this%w3, n, gs_op_add, &
382 this%gs_event)
383 call device_event_sync(this%gs_event)
384 call rotate_cyc(w1_d, w2_d, w3_d, 0, coef)
385
386 call bc_projector%apply(this%w1, this%w2, this%w3, n)
387
388 call device_vdot3(tmp_d, w1_d, w2_d, w3_d, p1_d, p2_d, p3_d, n)
389
390 pap = device_glsc2(tmp_d, coef%mult_d, n)
391
392 alpha = rtz1 / pap
393 alphm = -alpha
394 call device_opadd2cm(x%x_d, y%x_d, z%x_d, &
395 p1_d, p2_d, p3_d, alpha, n, gdim)
396 call device_opadd2cm(r1_d, r2_d, r3_d, &
397 w1_d, w2_d, w3_d, alphm, n, gdim)
398 call device_vdot3(tmp_d, r1_d, r2_d, r3_d, r1_d, r2_d, r3_d, n)
399
400 rtr = device_glsc2(tmp_d, coef%mult_d, n)
401 if (iter .eq. 1) rtr0 = rtr
402 rnorm = sqrt(rtr) * norm_fac
403 call this%monitor_iter(iter, rnorm)
404 if (rnorm .lt. this%abs_tol) then
405 exit
406 end if
407 end do
408 end associate
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 function cg_cpld_device_solve
415
416end module cg_cpld_device
__device__ T solve(const T u, const T y, const T guess, const T nu, const T kappa, const T B)
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
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Apply cyclic boundary condition to a vector field.
Defines a Matrix-vector product.
Definition ax.f90:34
Defines a coupled Conjugate Gradient methods for accelerators.
subroutine cg_cpld_device_free(this)
Deallocate a device based PCG solver.
type(ksp_monitor_t) function, dimension(3) cg_cpld_device_solve(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard PCG solve.
subroutine cg_cpld_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device based PCG solver.
type(ksp_monitor_t) function cg_cpld_device_nop(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Coefficients.
Definition coef.f90:34
subroutine, public device_add2s1(a_d, b_d, c1, n, strm)
subroutine, public device_rzero(a_d, n, strm)
Zero a real vector.
subroutine, public device_vdot3(dot_d, u1_d, u2_d, u3_d, v1_d, v2_d, v3_d, n, strm)
Compute a dot product (3-d version) assuming vector components etc.
subroutine, public device_copy(a_d, b_d, n, strm)
Copy a vector .
real(kind=rp) function, public device_glsc2(a_d, b_d, n, strm)
Weighted inner product .
subroutine, public device_opadd2cm(a1_d, a2_d, a3_d, b1_d, b2_d, b3_d, c, n, gdim)
Device abstraction, common interface for various accelerators.
Definition device.F90:34
subroutine, public device_event_sync(event)
Synchronize an event.
Definition device.F90:1667
subroutine, public device_event_destroy(event)
Destroy a device event.
Definition device.F90:1623
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.
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
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Operators.
Definition operators.f90:34
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...
Base type for a matrix-vector product providing .
Definition ax.f90:43
Device based coupled preconditioned conjugate gradient method.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:93
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.