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