Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
pc_jacobi_device.F90
Go to the documentation of this file.
1! Copyright (c) 2021-2025, 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 precon, only : pc_t
36 use bc_list, only : bc_list_t
37 use json_module, only : json_file
38 use coefs, only : coef_t
39 use dofmap, only : dofmap_t
40 use num_types, only : rp
45 use gather_scatter, only : gs_t, gs_op_add
46 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
47 implicit none
48 private
49
51 type, public, extends(pc_t) :: device_jacobi_t
52 real(kind=rp), allocatable :: d(:,:,:,:)
53 type(c_ptr) :: d_d = c_null_ptr
54 type(gs_t), pointer :: gs_h
55 type(dofmap_t), pointer :: dof
56 type(coef_t), pointer :: coef
57 type(c_ptr) :: gs_event = c_null_ptr
58 contains
59 procedure, pass(this) :: init => device_jacobi_init
60 procedure, pass(this) :: init_from_components => &
62 procedure, pass(this) :: free => device_jacobi_free
63 procedure, pass(this) :: solve => device_jacobi_solve
64 procedure, pass(this) :: update => device_jacobi_update
65 end type device_jacobi_t
66
67 interface
68 subroutine hip_jacobi_update(d_d, dxt_d, dyt_d, dzt_d, &
69 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
70 bind(c, name='hip_jacobi_update')
71 use, intrinsic :: iso_c_binding
72 type(c_ptr), value :: d_d, dxt_d, dyt_d, dzt_d
73 type(c_ptr), value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
74 integer(c_int) :: nelv, lx
75 end subroutine hip_jacobi_update
76 end interface
77
78 interface
79 subroutine cuda_jacobi_update(d_d, dxt_d, dyt_d, dzt_d, &
80 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
81 bind(c, name='cuda_jacobi_update')
82 use, intrinsic :: iso_c_binding
83 type(c_ptr), value :: d_d, dxt_d, dyt_d, dzt_d
84 type(c_ptr), value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
85 integer(c_int) :: nelv, lx
86 end subroutine cuda_jacobi_update
87 end interface
88
89 interface
90 subroutine opencl_jacobi_update(d_d, dxt_d, dyt_d, dzt_d, &
91 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
92 bind(c, name='opencl_jacobi_update')
93 use, intrinsic :: iso_c_binding
94 type(c_ptr), value :: d_d, dxt_d, dyt_d, dzt_d
95 type(c_ptr), value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
96 integer(c_int) :: nelv, lx
97 end subroutine opencl_jacobi_update
98 end interface
99
100 interface
101 subroutine metal_jacobi_update(d_d, dxt_d, dyt_d, dzt_d, &
102 G11_d, G22_d, G33_d, G12_d, G13_d, G23_d, nelv, lx) &
103 bind(c, name='metal_jacobi_update')
104 use, intrinsic :: iso_c_binding
105 type(c_ptr), value :: d_d, dxt_d, dyt_d, dzt_d
106 type(c_ptr), value :: G11_d, G22_d, G33_d, G12_d, G13_d, G23_d
107 integer(c_int) :: nelv, lx
108 end subroutine metal_jacobi_update
109 end interface
110
111contains
112
115 subroutine device_jacobi_init(this, coef, bclst, json)
116 class(device_jacobi_t), intent(inout), target :: this
117 type(coef_t), intent(in), target :: coef
118 type(bc_list_t), intent(inout), target :: bclst
119 type(json_file), intent(inout) :: json
120
121 call this%init_from_components(coef, coef%dof, coef%gs_h)
122
123 end subroutine device_jacobi_init
124
125 subroutine device_jacobi_init_from_components(this, coef, dof, gs_h)
126 class(device_jacobi_t), intent(inout) :: this
127 type(coef_t), intent(in), target :: coef
128 type(dofmap_t), intent(in), target :: dof
129 type(gs_t), intent(inout), target :: gs_h
130
131 call this%free()
132
133 this%gs_h => gs_h
134 this%dof => dof
135 this%coef => coef
136
137 allocate(this%d(dof%Xh%lx,dof%Xh%ly,dof%Xh%lz, dof%msh%nelv))
138
139 call device_map(this%d, this%d_d, size(this%d))
140
141 call device_event_create(this%gs_event, 2)
142
143 call device_jacobi_update(this)
144
146
147 subroutine device_jacobi_free(this)
148 class(device_jacobi_t), intent(inout) :: this
149
150 if (allocated(this%d)) then
151 if (c_associated(this%d_d)) then
152 call device_unmap(this%d, this%d_d)
153 end if
154 deallocate(this%d)
155 end if
156
157 if (c_associated(this%gs_event)) then
158 call device_event_destroy(this%gs_event)
159 end if
160
161 nullify(this%dof)
162 nullify(this%gs_h)
163 nullify(this%coef)
164 end subroutine device_jacobi_free
165
168 subroutine device_jacobi_solve(this, z, r, n)
169 integer, intent(in) :: n
170 class(device_jacobi_t), intent(inout) :: this
171 real(kind=rp), dimension(n), intent(inout) :: z
172 real(kind=rp), dimension(n), intent(inout) :: r
173 type(c_ptr) :: z_d, r_d
174
175 z_d = device_get_ptr(z)
176 r_d = device_get_ptr(r)
177
178 call device_col3(z_d, r_d, this%d_d, n)
179
180 end subroutine device_jacobi_solve
181
182 subroutine device_jacobi_update(this)
183 class(device_jacobi_t), intent(inout) :: this
184 integer :: lz, ly, lx
185 associate(dof => this%dof, coef => this%coef, xh => this%dof%Xh, &
186 gs_h => this%gs_h, nelv => this%dof%msh%nelv)
187
188 lx = xh%lx
189 ly = xh%ly
190 lz = xh%lz
191
192#ifdef HAVE_HIP
193 call hip_jacobi_update(this%d_d, xh%dxt_d, xh%dyt_d, xh%dzt_d, &
194 coef%G11_d, coef%G22_d, coef%G33_d, &
195 coef%G12_d, coef%G13_d, coef%G23_d, &
196 nelv, lx)
197#elif HAVE_CUDA
198 call cuda_jacobi_update(this%d_d, xh%dxt_d, xh%dyt_d, xh%dzt_d, &
199 coef%G11_d, coef%G22_d, coef%G33_d, &
200 coef%G12_d, coef%G13_d, coef%G23_d, &
201 nelv, lx)
202#elif HAVE_OPENCL
203 call opencl_jacobi_update(this%d_d, xh%dxt_d, xh%dyt_d, xh%dzt_d, &
204 coef%G11_d, coef%G22_d, coef%G33_d, &
205 coef%G12_d, coef%G13_d, coef%G23_d, &
206 nelv, lx)
207#elif HAVE_METAL
208 call metal_jacobi_update(this%d_d, xh%dxt_d, xh%dyt_d, xh%dzt_d, &
209 coef%G11_d, coef%G22_d, coef%G33_d, &
210 coef%G12_d, coef%G13_d, coef%G23_d, &
211 nelv, lx)
212#endif
213
214 call device_col2(this%d_d, coef%h1_d, coef%dof%size())
215
216 if (coef%ifh2) then
217 call device_addcol3(this%d_d, coef%h2_d, coef%B_d, coef%dof%size())
218 end if
219
220 call gs_h%op(this%d, dof%size(), gs_op_add, this%gs_event)
221 call device_event_sync(this%gs_event)
222
223 call device_invcol1(this%d_d, dof%size())
224 end associate
225 end subroutine device_jacobi_update
226
227end module device_jacobi
__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 list of bc_t.
Definition bc_list.f90:34
Coefficients.
Definition coef.f90:34
Jacobi preconditioner accelerator backend.
subroutine device_jacobi_free(this)
subroutine device_jacobi_init_from_components(this, coef, dof, gs_h)
subroutine device_jacobi_update(this)
subroutine device_jacobi_init(this, coef, bclst, json)
Initialize from the case file. The boundary conditions and the JSON dictionary are not needed by this...
subroutine device_jacobi_solve(this, z, r, n)
The jacobi preconditioner where .
subroutine, public device_addcol3(a_d, b_d, c_d, n, strm)
Returns .
subroutine, public device_invcol1(a_d, n, strm)
Invert a vector .
subroutine, public device_col2(a_d, b_d, n, strm)
Vector multiplication .
subroutine, public device_col3(a_d, b_d, c_d, n, strm)
Vector multiplication with 3 vectors .
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 mapping of the degrees of freedom.
Definition dofmap.f90:35
Gather-scatter.
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Krylov preconditioner.
Definition precon.f90:34
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:136
Defines a jacobi preconditioner.
Gather-scatter kernel.
Defines a canonical Krylov preconditioner.
Definition precon.f90:44