Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
cg_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!
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
45 use math, only : abscmp
50 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
51 implicit none
52 private
53
55 type, public, extends(ksp_t) :: cg_device_t
56 real(kind=rp), allocatable :: w(:)
57 real(kind=rp), allocatable :: r(:)
58 real(kind=rp), allocatable :: p(:)
59 real(kind=rp), allocatable :: z(:)
60 type(c_ptr) :: w_d = c_null_ptr
61 type(c_ptr) :: r_d = c_null_ptr
62 type(c_ptr) :: p_d = c_null_ptr
63 type(c_ptr) :: z_d = c_null_ptr
64 type(c_ptr) :: gs_event = c_null_ptr
65 contains
66 procedure, pass(this) :: init => cg_device_init
67 procedure, pass(this) :: free => cg_device_free
68 procedure, pass(this) :: solve => cg_device_solve
69 procedure, pass(this) :: solve_coupled => cg_device_solve_coupled
70 end type cg_device_t
71
72contains
73
75 subroutine cg_device_init(this, n, max_iter, M, rel_tol, abs_tol, monitor)
76 class(cg_device_t), target, intent(inout) :: this
77 class(pc_t), optional, intent(in), target :: M
78 integer, intent(in) :: n
79 integer, intent(in) :: max_iter
80 real(kind=rp), optional, intent(in) :: rel_tol
81 real(kind=rp), optional, intent(in) :: abs_tol
82 logical, optional, intent(in) :: monitor
83
84 call this%free()
85
86 allocate(this%w(n))
87 allocate(this%r(n))
88 allocate(this%p(n))
89 allocate(this%z(n))
90
91 call device_map(this%z, this%z_d, n)
92 call device_map(this%p, this%p_d, n)
93 call device_map(this%r, this%r_d, n)
94 call device_map(this%w, this%w_d, n)
95
96 if (present(m)) then
97 this%M => m
98 end if
99
100 if (present(rel_tol) .and. present(abs_tol) .and. present(monitor)) then
101 call this%ksp_init(max_iter, rel_tol, abs_tol, monitor = monitor)
102 else if (present(rel_tol) .and. present(abs_tol)) then
103 call this%ksp_init(max_iter, rel_tol, abs_tol)
104 else if (present(monitor) .and. present(abs_tol)) then
105 call this%ksp_init(max_iter, abs_tol = abs_tol, monitor = monitor)
106 else if (present(rel_tol) .and. present(monitor)) then
107 call this%ksp_init(max_iter, rel_tol, monitor = monitor)
108 else if (present(rel_tol)) then
109 call this%ksp_init(max_iter, rel_tol = rel_tol)
110 else if (present(abs_tol)) then
111 call this%ksp_init(max_iter, abs_tol = abs_tol)
112 else if (present(monitor)) then
113 call this%ksp_init(max_iter, monitor = monitor)
114 else
115 call this%ksp_init(max_iter)
116 end if
117
118 call device_event_create(this%gs_event, 2)
119 end subroutine cg_device_init
120
122 subroutine cg_device_free(this)
123 class(cg_device_t), intent(inout) :: this
124
125 call this%ksp_free()
126
127 if (allocated(this%w)) then
128 if (c_associated(this%w_d)) then
129 call device_unmap(this%w, this%w_d)
130 end if
131 deallocate(this%w)
132 end if
133
134 if (allocated(this%r)) then
135 if (c_associated(this%r_d)) then
136 call device_unmap(this%r, this%r_d)
137 end if
138 deallocate(this%r)
139 end if
140
141 if (allocated(this%p)) then
142 if (c_associated(this%p_d)) then
143 call device_unmap(this%p, this%p_d)
144 end if
145 deallocate(this%p)
146 end if
147
148 if (allocated(this%z)) then
149 if (c_associated(this%z_d)) then
150 call device_unmap(this%z, this%z_d)
151 end if
152 deallocate(this%z)
153 end if
154
155 nullify(this%M)
156
157 if (c_associated(this%gs_event)) then
158 call device_event_destroy(this%gs_event)
159 end if
160
161 end subroutine cg_device_free
162
164 function cg_device_solve(this, Ax, x, f, n, coef, bc_projector, gs_h, niter) &
165 result(ksp_results)
166 class(cg_device_t), intent(inout) :: this
167 class(ax_t), intent(in) :: ax
168 type(field_t), intent(inout) :: x
169 integer, intent(in) :: n
170 real(kind=rp), dimension(n), intent(in) :: f
171 type(coef_t), intent(inout) :: coef
172 class(scalar_bc_projector_t), intent(inout) :: bc_projector
173 type(gs_t), intent(inout) :: gs_h
174 type(ksp_monitor_t) :: ksp_results
175 integer, optional, intent(in) :: niter
176 real(kind=rp), parameter :: one = 1.0
177 real(kind=rp), parameter :: zero = 0.0
178 integer :: iter, max_iter
179 real(kind=rp) :: rnorm, rtr, rtr0, rtz2, rtz1
180 real(kind=rp) :: beta, pap, alpha, alphm, norm_fac
181 type(c_ptr) :: f_d
182
183 f_d = device_get_ptr(f)
184
185 if (present(niter)) then
186 max_iter = niter
187 else
188 max_iter = this%max_iter
189 end if
190 norm_fac = one/sqrt(coef%volume)
191
192 rtz1 = one
193 call device_rzero(x%x_d, n)
194 call device_rzero(this%p_d, n)
195 call device_copy(this%r_d, f_d, n)
196
197 rtr = device_glsc3(this%r_d, coef%mult_d, this%r_d, n)
198 rnorm = sqrt(rtr)*norm_fac
199 ksp_results%res_start = rnorm
200 ksp_results%res_final = rnorm
201 ksp_results%iter = 0
202 if (abscmp(rnorm, zero)) then
203 ksp_results%converged = .true.
204 return
205 end if
206 call this%monitor_start('CG')
207 do iter = 1, max_iter
208 call this%M%solve(this%z, this%r, n)
209 rtz2 = rtz1
210 rtz1 = device_glsc3(this%r_d, coef%mult_d, this%z_d, n)
211 beta = rtz1 / rtz2
212 if (iter .eq. 1) beta = zero
213 call device_add2s1(this%p_d, this%z_d, beta, n)
214
215 call ax%compute(this%w, this%p, coef, x%msh, x%Xh)
216 call gs_h%op(this%w, n, gs_op_add, this%gs_event)
217 call device_event_sync(this%gs_event)
218 call bc_projector%apply(this%w, n)
219
220 pap = device_glsc3(this%w_d, coef%mult_d, this%p_d, n)
221
222 alpha = rtz1 / pap
223 alphm = -alpha
224 call device_add2s2(x%x_d, this%p_d, alpha, n)
225 call device_add2s2(this%r_d, this%w_d, alphm, n)
226
227 rtr = device_glsc3(this%r_d, coef%mult_d, this%r_d, n)
228 if (iter .eq. 1) rtr0 = rtr
229 rnorm = sqrt(rtr)*norm_fac
230 call this%monitor_iter(iter, rnorm)
231 if (rnorm .lt. this%abs_tol) then
232 exit
233 end if
234 end do
235 call this%monitor_stop()
236 ksp_results%res_final = rnorm
237 ksp_results%iter = iter
238 ksp_results%converged = this%is_converged(iter, rnorm)
239
240 end function cg_device_solve
241
243 function cg_device_solve_coupled(this, Ax, x, y, z, fx, fy, fz, &
244 n, coef, bc_projector, gs_h, niter) result(ksp_results)
245 class(cg_device_t), intent(inout) :: this
246 class(ax_t), intent(in) :: ax
247 type(field_t), intent(inout) :: x
248 type(field_t), intent(inout) :: y
249 type(field_t), intent(inout) :: z
250 integer, intent(in) :: n
251 real(kind=rp), dimension(n), intent(in) :: fx
252 real(kind=rp), dimension(n), intent(in) :: fy
253 real(kind=rp), dimension(n), intent(in) :: fz
254 type(coef_t), intent(inout) :: coef
255 class(vector_bc_projector_t), intent(inout) :: bc_projector
256 type(gs_t), intent(inout) :: gs_h
257 type(ksp_monitor_t), dimension(3) :: ksp_results
258 integer, optional, intent(in) :: niter
259 type(scalar_bc_projector_t), pointer :: bc_x, bc_y, bc_z
260
261 call vector_bc_projector_components(bc_projector, bc_x, bc_y, bc_z)
262 ksp_results(1) = this%solve(ax, x, fx, n, coef, bc_x, gs_h, niter)
263 ksp_results(2) = this%solve(ax, y, fy, n, coef, bc_y, gs_h, niter)
264 ksp_results(3) = this%solve(ax, z, fz, n, coef, bc_z, gs_h, niter)
265
266 end function cg_device_solve_coupled
267
268end module cg_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
Defines a Matrix-vector product.
Definition ax.f90:34
Defines various Conjugate Gradient methods for accelerators.
Definition cg_device.f90:34
type(ksp_monitor_t) function, dimension(3) cg_device_solve_coupled(this, ax, x, y, z, fx, fy, fz, n, coef, bc_projector, gs_h, niter)
Standard PCG coupled solve.
subroutine cg_device_init(this, n, max_iter, m, rel_tol, abs_tol, monitor)
Initialise a device based PCG solver.
Definition cg_device.f90:76
subroutine cg_device_free(this)
Deallocate a device based PCG solver.
type(ksp_monitor_t) function cg_device_solve(this, ax, x, f, n, coef, bc_projector, gs_h, niter)
Standard PCG solve.
Coefficients.
Definition coef.f90:34
subroutine, public device_add2s1(a_d, b_d, c1, 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_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
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
Krylov preconditioner.
Definition precon.f90:34
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 .
Definition ax.f90:43
Device based preconditioned conjugate gradient method.
Definition cg_device.f90:55
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.