Neko 1.99.6
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
42 use bc_list, only : bc_list_t
43 use math, only : abscmp
49 use utils, only : neko_error
50 use operators, only : rotate_cyc
51 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
52 implicit none
53 private
54
56 type, public, extends(ksp_t) :: cg_cpld_device_t
57 real(kind=rp), allocatable :: w1(:)
58 real(kind=rp), allocatable :: w2(:)
59 real(kind=rp), allocatable :: w3(:)
60 real(kind=rp), allocatable :: r1(:)
61 real(kind=rp), allocatable :: r2(:)
62 real(kind=rp), allocatable :: r3(:)
63 real(kind=rp), allocatable :: p1(:)
64 real(kind=rp), allocatable :: p2(:)
65 real(kind=rp), allocatable :: p3(:)
66 real(kind=rp), allocatable :: z1(:)
67 real(kind=rp), allocatable :: z2(:)
68 real(kind=rp), allocatable :: z3(:)
69 real(kind=rp), allocatable :: tmp(:)
70
71
72 type(c_ptr) :: w1_d = c_null_ptr
73 type(c_ptr) :: w2_d = c_null_ptr
74 type(c_ptr) :: w3_d = c_null_ptr
75
76 type(c_ptr) :: r1_d = c_null_ptr
77 type(c_ptr) :: r2_d = c_null_ptr
78 type(c_ptr) :: r3_d = c_null_ptr
79
80 type(c_ptr) :: p1_d = c_null_ptr
81 type(c_ptr) :: p2_d = c_null_ptr
82 type(c_ptr) :: p3_d = c_null_ptr
83
84 type(c_ptr) :: z1_d = c_null_ptr
85 type(c_ptr) :: z2_d = c_null_ptr
86 type(c_ptr) :: z3_d = c_null_ptr
87
88 type(c_ptr) :: tmp_d = c_null_ptr
89
90 type(c_ptr) :: gs_event = c_null_ptr
91 contains
92 procedure, pass(this) :: init => cg_cpld_device_init
93 procedure, pass(this) :: free => cg_cpld_device_free
94 procedure, pass(this) :: solve => cg_cpld_device_nop
95 procedure, pass(this) :: solve_coupled => cg_cpld_device_solve
96 end type cg_cpld_device_t
97
98contains
99
101 subroutine cg_cpld_device_init(this, n, max_iter, M, rel_tol, abs_tol, &
102 monitor)
103 class(cg_cpld_device_t), target, intent(inout) :: this
104 class(pc_t), optional, intent(in), target :: M
105 integer, intent(in) :: n
106 integer, intent(in) :: max_iter
107 real(kind=rp), optional, intent(in) :: rel_tol
108 real(kind=rp), optional, intent(in) :: abs_tol
109 logical, optional, intent(in) :: monitor
110
111 call this%free()
112
113 allocate(this%w1(n))
114 allocate(this%w2(n))
115 allocate(this%w3(n))
116 allocate(this%r1(n))
117 allocate(this%r2(n))
118 allocate(this%r3(n))
119 allocate(this%p1(n))
120 allocate(this%p2(n))
121 allocate(this%p3(n))
122 allocate(this%z1(n))
123 allocate(this%z2(n))
124 allocate(this%z3(n))
125 allocate(this%tmp(n))
126
127 call device_map(this%tmp, this%tmp_d, n)
128 call device_map(this%z1, this%z1_d, n)
129 call device_map(this%z2, this%z2_d, n)
130 call device_map(this%z3, this%z3_d, n)
131 call device_map(this%p1, this%p1_d, n)
132 call device_map(this%p2, this%p2_d, n)
133 call device_map(this%p3, this%p3_d, n)
134 call device_map(this%r1, this%r1_d, n)
135 call device_map(this%r2, this%r2_d, n)
136 call device_map(this%r3, this%r3_d, n)
137 call device_map(this%w1, this%w1_d, n)
138 call device_map(this%w2, this%w2_d, n)
139 call device_map(this%w3, this%w3_d, n)
140
141 if (present(m)) then
142 this%M => m
143 end if
144
145 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
146 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
147 else if (present(rel_tol) .and. present(abs_tol)) then
148 call this%ksp_init(max_iter, rel_tol, abs_tol)
149 else if (present(monitor) .and. present(abs_tol)) then
150 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
151 else if (present(rel_tol) .and. present(monitor)) then
152 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
153 else if (present(rel_tol)) then
154 call this%ksp_init(max_iter, rel_tol = rel_tol)
155 else if (present(abs_tol)) then
156 call this%ksp_init(max_iter, abs_tol = abs_tol)
157 else if (present(monitor)) then
158 call this%ksp_init(max_iter, monitor = monitor)
159 else
160 call this%ksp_init(max_iter)
161 end if
162
163 call device_event_create(this%gs_event, 2)
164 end subroutine cg_cpld_device_init
165
167 subroutine cg_cpld_device_free(this)
168 class(cg_cpld_device_t), intent(inout) :: this
169
170 call this%ksp_free()
171
172 if (allocated(this%w1)) then
173 if (c_associated(this%w1_d)) then
174 call device_unmap(this%w1, this%w1_d)
175 end if
176 deallocate(this%w1)
177 end if
178
179 if (allocated(this%w2)) then
180 if (c_associated(this%w2_d)) then
181 call device_unmap(this%w2, this%w2_d)
182 end if
183 deallocate(this%w2)
184 end if
185
186 if (allocated(this%w3)) then
187 if (c_associated(this%w3_d)) then
188 call device_unmap(this%w3, this%w3_d)
189 end if
190 deallocate(this%w3)
191 end if
192
193 if (allocated(this%r1)) then
194 if (c_associated(this%r1_d)) then
195 call device_unmap(this%r1, this%r1_d)
196 end if
197 deallocate(this%r1)
198 end if
199
200 if (allocated(this%r2)) then
201 if (c_associated(this%r2_d)) then
202 call device_unmap(this%r2, this%r2_d)
203 end if
204 deallocate(this%r2)
205 end if
206
207 if (allocated(this%r3)) then
208 if (c_associated(this%r3_d)) then
209 call device_unmap(this%r3, this%r3_d)
210 end if
211 deallocate(this%r3)
212 end if
213
214 if (allocated(this%p1)) then
215 if (c_associated(this%p1_d)) then
216 call device_unmap(this%p1, this%p1_d)
217 end if
218 deallocate(this%p1)
219 end if
220
221 if (allocated(this%p2)) then
222 if (c_associated(this%p2_d)) then
223 call device_unmap(this%p2, this%p2_d)
224 end if
225 deallocate(this%p2)
226 end if
227
228 if (allocated(this%p3)) then
229 if (c_associated(this%p3_d)) then
230 call device_unmap(this%p3, this%p3_d)
231 end if
232 deallocate(this%p3)
233 end if
234
235 if (allocated(this%z1)) then
236 if (c_associated(this%z1_d)) then
237 call device_unmap(this%z1, this%z1_d)
238 end if
239 deallocate(this%z1)
240 end if
241
242 if (allocated(this%z2)) then
243 if (c_associated(this%z2_d)) then
244 call device_unmap(this%z2, this%z2_d)
245 end if
246 deallocate(this%z2)
247 end if
248
249 if (allocated(this%z3)) then
250 if (c_associated(this%z3_d)) then
251 call device_unmap(this%z3, this%z3_d)
252 end if
253 deallocate(this%z3)
254 end if
255
256 if (allocated(this%tmp)) then
257 if (c_associated(this%tmp_d)) then
258 call device_unmap(this%tmp, this%tmp_d)
259 end if
260 deallocate(this%tmp)
261 end if
262
263 nullify(this%M)
264
265 if (c_associated(this%gs_event)) then
266 call device_event_destroy(this%gs_event)
267 end if
268
269 end subroutine cg_cpld_device_free
270
271 function cg_cpld_device_nop(this, Ax, x, f, n, coef, blst, gs_h, niter) &
272 result(ksp_results)
273 class(cg_cpld_device_t), intent(inout) :: this
274 class(ax_t), intent(in) :: ax
275 type(field_t), intent(inout) :: x
276 integer, intent(in) :: n
277 real(kind=rp), dimension(n), intent(in) :: f
278 type(coef_t), intent(inout) :: coef
279 type(bc_list_t), intent(inout) :: blst
280 type(gs_t), intent(inout) :: gs_h
281 type(ksp_monitor_t) :: ksp_results
282 integer, optional, intent(in) :: niter
283
284 ! Throw and error
285 call neko_error('The cpldcg solver is only defined for coupled solves')
286
287 ksp_results%res_final = 0.0
288 ksp_results%iter = 0
289 end function cg_cpld_device_nop
290
292 function cg_cpld_device_solve(this, Ax, x, y, z, fx, fy, fz, &
293 n, coef, blstx, blsty, blstz, gs_h, niter) result(ksp_results)
294 class(cg_cpld_device_t), intent(inout) :: this
295 class(ax_t), intent(in) :: ax
296 type(field_t), intent(inout) :: x
297 type(field_t), intent(inout) :: y
298 type(field_t), intent(inout) :: z
299 integer, intent(in) :: n
300 real(kind=rp), dimension(n), intent(in) :: fx
301 real(kind=rp), dimension(n), intent(in) :: fy
302 real(kind=rp), dimension(n), intent(in) :: fz
303 type(coef_t), intent(inout) :: coef
304 type(bc_list_t), intent(inout) :: blstx
305 type(bc_list_t), intent(inout) :: blsty
306 type(bc_list_t), intent(inout) :: blstz
307 type(gs_t), intent(inout) :: gs_h
308 type(ksp_monitor_t), dimension(3) :: ksp_results
309 integer, optional, intent(in) :: niter
310 integer :: i, iter, max_iter
311 real(kind=rp) :: rnorm, rtr, rtr0, rtz2, rtz1
312 real(kind=rp) :: beta, pap, alpha, alphm, norm_fac
313 integer, parameter :: gdim = 3
314 type(c_ptr) :: fx_d
315 type(c_ptr) :: fy_d
316 type(c_ptr) :: fz_d
317
318 fx_d = device_get_ptr(fx)
319 fy_d = device_get_ptr(fy)
320 fz_d = device_get_ptr(fz)
321
322 if (present(niter)) then
323 max_iter = niter
324 else
325 max_iter = this%max_iter
326 end if
327 norm_fac = 1.0_rp / sqrt(coef%volume)
328
329 associate(p1_d => this%p1_d, p2_d => this%p2_d, p3_d => this%p3_d, &
330 z1_d => this%z1_d, z2_d => this%z2_d, z3_d => this%z3_d, &
331 r1_d => this%r1_d, r2_d => this%r2_d, r3_d => this%r3_d, &
332 w1_d => this%w1_d, w2_d => this%w2_d, w3_d => this%w3_d, &
333 tmp_d => this%tmp_d)
334
335 rtz1 = 1.0_rp
336 call device_rzero(x%x_d, n)
337 call device_rzero(y%x_d, n)
338 call device_rzero(z%x_d, n)
339 call device_rzero(p1_d, n)
340 call device_rzero(p2_d, n)
341 call device_rzero(p3_d, n)
342 call device_rzero(z1_d, n)
343 call device_rzero(z2_d, n)
344 call device_rzero(z3_d, n)
345 call device_copy(r1_d, fx_d, n)
346 call device_copy(r2_d, fy_d, n)
347 call device_copy(r3_d, fz_d, n)
348 call device_vdot3(tmp_d, r1_d, r2_d, r3_d, r1_d, r2_d, r3_d, n)
349
350
351 rtr = device_glsc2(tmp_d, coef%mult_d, n)
352 rnorm = sqrt(rtr)*norm_fac
353 ksp_results%res_start = rnorm
354 ksp_results%res_final = rnorm
355 ksp_results%iter = 0
356 if (abscmp(rnorm, 0.0_rp)) then
357 ksp_results%converged = .true.
358 return
359 end if
360
361 call this%monitor_start('device_cpldCG')
362 do iter = 1, max_iter
363 call this%M%solve(this%z1, this%r1, n)
364 call this%M%solve(this%z2, this%r2, n)
365 call this%M%solve(this%z3, this%r3, n)
366 rtz2 = rtz1
367
368 call device_vdot3(tmp_d, z1_d, z2_d, z3_d, r1_d, r2_d, r3_d, n)
369
370 rtz1 = device_glsc2(tmp_d, coef%mult_d, n)
371
372 beta = rtz1 / rtz2
373 if (iter .eq. 1) beta = 0.0_rp
374 call device_add2s1(p1_d, z1_d, beta, n)
375 call device_add2s1(p2_d, z2_d, beta, n)
376 call device_add2s1(p3_d, z3_d, beta, n)
377
378 call ax%compute_vector(this%w1, this%w2, this%w3, &
379 this%p1, this%p2, this%p3, coef, x%msh, x%Xh)
380
381 call rotate_cyc(w1_d, w2_d, w3_d, 1, coef)
382 call gs_h%op(this%w1, this%w2, this%w3, n, gs_op_add, &
383 this%gs_event)
384 call device_event_sync(this%gs_event)
385 call rotate_cyc(w1_d, w2_d, w3_d, 0, coef)
386
387 call blstx%apply(this%w1, n)
388 call blsty%apply(this%w2, n)
389 call blstz%apply(this%w3, n)
390
391 call device_vdot3(tmp_d, w1_d, w2_d, w3_d, p1_d, p2_d, p3_d, n)
392
393 pap = device_glsc2(tmp_d, coef%mult_d, n)
394
395 alpha = rtz1 / pap
396 alphm = -alpha
397 call device_opadd2cm(x%x_d, y%x_d, z%x_d, &
398 p1_d, p2_d, p3_d, alpha, n, gdim)
399 call device_opadd2cm(r1_d, r2_d, r3_d, &
400 w1_d, w2_d, w3_d, alphm, n, gdim)
401 call device_vdot3(tmp_d, r1_d, r2_d, r3_d, r1_d, r2_d, r3_d, n)
402
403 rtr = device_glsc2(tmp_d, coef%mult_d, n)
404 if (iter .eq. 1) rtr0 = rtr
405 rnorm = sqrt(rtr) * norm_fac
406 call this%monitor_iter(iter, rnorm)
407 if (rnorm .lt. this%abs_tol) then
408 exit
409 end if
410 end do
411 end associate
412 call this%monitor_stop()
413 ksp_results%res_final = rnorm
414 ksp_results%iter = iter
415 ksp_results%converged = this%is_converged(iter, rnorm)
416
417 end function cg_cpld_device_solve
418
419end 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 list of bc_t.
Definition bc_list.f90:34
Defines a coupled Conjugate Gradient methods for accelerators.
type(ksp_monitor_t) function, dimension(3) cg_cpld_device_solve(this, ax, x, y, z, fx, fy, fz, n, coef, blstx, blsty, blstz, gs_h, niter)
Standard PCG solve.
type(ksp_monitor_t) function cg_cpld_device_nop(this, ax, x, f, n, coef, blst, gs_h, niter)
subroutine cg_cpld_device_free(this)
Deallocate a device based PCG solver.
subroutine cg_cpld_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device based PCG solver.
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:51
Definition math.f90:60
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
Operators.
Definition operators.f90:34
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
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:63
Gather-scatter kernel.
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