Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
facet_normal.f90
Go to the documentation of this file.
1! Copyright (c) 2020-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 num_types, only : rp
37 use math, only : cfill_mask
40 use vector, only : vector_t
41 use coefs, only : coef_t
42 use bc, only : bc_t, bc_dirichlet
44 use json_module, only : json_file
45 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
46 use htable, only : htable_i4_t
49 use time_state, only : time_state_t
50 implicit none
51 private
52
54 type, public, extends(bc_t) :: facet_normal_t
55 integer, allocatable :: unique_mask(:)
56 integer, allocatable :: msk_to_unique(:)
57 type(c_ptr) :: unique_mask_d = c_null_ptr
58 type(vector_t) :: nx, ny, nz, work
59 contains
60 procedure, pass(this) :: apply_scalar => facet_normal_apply_scalar
61 procedure, pass(this) :: apply_scalar_dev => facet_normal_apply_scalar_dev
62 procedure, pass(this) :: apply_vector => facet_normal_apply_vector
63 procedure, pass(this) :: apply_vector_dev => facet_normal_apply_vector_dev
64 procedure, pass(this) :: apply_surfvec => facet_normal_apply_surfvec
65 procedure, pass(this) :: apply_surfvec_dev => &
67 ! > Recompute normals
68 procedure, pass(this) :: recompute_normals => facet_normal_recompute_normals
70 procedure, pass(this) :: init => facet_normal_init
72 procedure, pass(this) :: init_from_components => &
75 procedure, pass(this) :: free => facet_normal_free
77 procedure, pass(this) :: finalize => facet_normal_finalize
78 end type facet_normal_t
79
80contains
81
85 subroutine facet_normal_init(this, coef, json)
86 class(facet_normal_t), intent(inout), target :: this
87 type(coef_t), target, intent(in) :: coef
88 type(json_file), intent(inout) :: json
89
90 call this%init_from_components(coef)
91 end subroutine facet_normal_init
92
95 subroutine facet_normal_init_from_components(this, coef)
96 class(facet_normal_t), intent(inout), target :: this
97 type(coef_t), target, intent(in) :: coef
98
99 call this%init_base(coef)
100 this%bc_type = bc_dirichlet
102
104 subroutine facet_normal_apply_scalar(this, x, n, time, strong)
105 class(facet_normal_t), intent(inout) :: this
106 integer, intent(in) :: n
107 real(kind=rp), intent(inout), dimension(n) :: x
108 type(time_state_t), intent(in), optional :: time
109 logical, intent(in), optional :: strong
110 end subroutine facet_normal_apply_scalar
111
113 subroutine facet_normal_apply_scalar_dev(this, x_d, time, strong, strm)
114 class(facet_normal_t), intent(inout), target :: this
115 type(c_ptr), intent(inout) :: x_d
116 type(time_state_t), intent(in), optional :: time
117 logical, intent(in), optional :: strong
118 type(c_ptr), intent(inout) :: strm
119
120 end subroutine facet_normal_apply_scalar_dev
121
123 subroutine facet_normal_apply_vector_dev(this, x_d, y_d, z_d, time, &
124 strong, strm)
125 class(facet_normal_t), intent(inout), target :: this
126 type(c_ptr), intent(inout) :: x_d
127 type(c_ptr), intent(inout) :: y_d
128 type(c_ptr), intent(inout) :: z_d
129 type(time_state_t), intent(in), optional :: time
130 logical, intent(in), optional :: strong
131 type(c_ptr), intent(inout) :: strm
132
133 end subroutine facet_normal_apply_vector_dev
134
136 subroutine facet_normal_apply_vector(this, x, y, z, n, time, strong)
137 class(facet_normal_t), intent(inout) :: this
138 integer, intent(in) :: n
139 real(kind=rp), intent(inout), dimension(n) :: x
140 real(kind=rp), intent(inout), dimension(n) :: y
141 real(kind=rp), intent(inout), dimension(n) :: z
142 type(time_state_t), intent(in), optional :: time
143 logical, intent(in), optional :: strong
144 end subroutine facet_normal_apply_vector
145
147 subroutine facet_normal_apply_surfvec(this, x, y, z, u, v, w, n, time)
148 class(facet_normal_t), intent(in) :: this
149 integer, intent(in) :: n
150 real(kind=rp), intent(inout), dimension(n) :: x
151 real(kind=rp), intent(inout), dimension(n) :: y
152 real(kind=rp), intent(inout), dimension(n) :: z
153 real(kind=rp), intent(inout), dimension(n) :: u
154 real(kind=rp), intent(inout), dimension(n) :: v
155 real(kind=rp), intent(inout), dimension(n) :: w
156 type(time_state_t), intent(in), optional :: time
157 integer :: i, m, k, idx(4), facet
158 real(kind=rp) :: normal(3), area
159
160 m = this%unique_mask(0)
161 ! Since apply_surfvec is called outside of the parallel region, we
162 ! need to open a separate parallel region here
163 !$omp parallel do
164 do i = 1, m
165 k = this%unique_mask(i)
166 x(k) = u(k) * this%nx%x(i)
167 y(k) = v(k) * this%ny%x(i)
168 z(k) = w(k) * this%nz%x(i)
169 end do
170 !$omp end parallel do
171
172 end subroutine facet_normal_apply_surfvec
173
175 subroutine facet_normal_apply_surfvec_dev(this, x_d, y_d, z_d, &
176 u_d, v_d, w_d, time, strm)
177 class(facet_normal_t), intent(in), target :: this
178 type(c_ptr) :: x_d, y_d, z_d, u_d, v_d, w_d
179 type(time_state_t), intent(in), optional :: time
180 type(c_ptr), optional :: strm
181 type(c_ptr) :: strm_
182 integer :: n, m
183
184 n = this%coef%dof%size()
185 m = this%unique_mask(0)
186
187 if (present(strm)) then
188 strm_ = strm
189 else
190 strm_ = glb_cmd_queue
191 end if
192
193 if (m .gt. 0) then
194 call device_masked_gather_copy_0(this%work%x_d, u_d, &
195 this%unique_mask_d, n, m, strm_)
196 call device_col2(this%work%x_d, this%nx%x_d, m, strm_)
197 call device_masked_scatter_copy_0(x_d, this%work%x_d, &
198 this%unique_mask_d, n, m, strm_)
199 call device_masked_gather_copy_0(this%work%x_d, v_d, &
200 this%unique_mask_d, n, m, strm_)
201 call device_col2(this%work%x_d, this%ny%x_d, m, strm_)
202 call device_masked_scatter_copy_0(y_d, this%work%x_d, &
203 this%unique_mask_d, n, m, strm_)
204 call device_masked_gather_copy_0(this%work%x_d, w_d, &
205 this%unique_mask_d, n, m, strm_)
206 call device_col2(this%work%x_d, this%nz%x_d, m, strm)
207 call device_masked_scatter_copy_0(z_d, this%work%x_d, &
208 this%unique_mask_d, n, m, strm_)
209 end if
210
211 end subroutine facet_normal_apply_surfvec_dev
212
214 subroutine facet_normal_free(this)
215 class(facet_normal_t), target, intent(inout) :: this
216
217 call this%free_base()
218 if (allocated(this%unique_mask)) then
219 if (neko_bcknd_device .eq. 1) then
220 call device_unmap(this%unique_mask, this%unique_mask_d)
221 end if
222 deallocate(this%unique_mask)
223 end if
224 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
225
226 call this%nx%free()
227 call this%ny%free()
228 call this%nz%free()
229 call this%work%free()
230
231 end subroutine facet_normal_free
232
234 subroutine facet_normal_finalize(this)
235 class(facet_normal_t), target, intent(inout) :: this
236 type(htable_i4_t) :: unique_point_idx
237 integer :: htable_data, rcode, i, j, idx(4), facet
238 real(kind=rp) :: area, normal(3)
239
240 ! Here and in recompute_normals(), which only runs after this
241 call this%coef%require_facets('facet_normal')
242
243 call this%finalize_base()
244
245 ! This part is purely needed to ensure that contributions
246 ! for all faces a point is on is properly summed up.
247 ! If one simply uses the original mask, if a point is on a corner
248 ! where both faces are on the boundary
249 ! one will only get the contribution from one face, not both
250 ! We solve this by adding up the normals of both faces for these points
251 ! and storing this sum in this%nx, this%ny, this%nz.
252 ! As both contributions are added already,
253 ! we also ensure that we only visit each point once
254 ! and create a new mask with only unique points (this%unique_mask).
255 if (allocated(this%unique_mask)) then
256 if (neko_bcknd_device .eq. 1) then
257 call device_unmap(this%unique_mask, this%unique_mask_d)
258 end if
259 deallocate(this%unique_mask)
260 end if
261 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
262
263 call unique_point_idx%init(this%facet_node_msk(0), htable_data)
264 j = 0
265 do i = 1, this%facet_node_msk(0)
266 if (unique_point_idx%get(this%facet_node_msk(i), &
267 htable_data) .ne. 0) then
268 j = j + 1
269 htable_data = j
270 call unique_point_idx%set(this%facet_node_msk(i), j)
271 end if
272 end do
273
274 ! Only allocate work vectors if size is non-zero
275 if (unique_point_idx%num_entries() .gt. 0 ) then
276 call this%nx%init(unique_point_idx%num_entries())
277 call this%ny%init(unique_point_idx%num_entries())
278 call this%nz%init(unique_point_idx%num_entries())
279 call this%work%init(unique_point_idx%num_entries())
280 end if
281 allocate(this%unique_mask(0:unique_point_idx%num_entries()))
282 allocate(this%msk_to_unique(this%facet_node_msk(0)))
283
284 this%unique_mask(0) = unique_point_idx%num_entries()
285 do i = 1, this%unique_mask(0)
286 this%unique_mask(i) = 0
287 end do
288
289
290 do i = 1, this%facet_node_msk(0)
291 rcode = unique_point_idx%get(this%facet_node_msk(i), htable_data)
292 if (rcode .ne. 0) call neko_error("Facet normal: htable get failed.")
293 this%unique_mask(htable_data) = this%facet_node_msk(i)
294
295 ! Save the slot so recompute_normals can use it without the hash table.
296 this%msk_to_unique(i) = htable_data
297 facet = this%facet(i)
298
299 idx = nonlinear_index(this%facet_node_msk(i), this%Xh%lx, this%Xh%lx, &
300 this%Xh%lx)
301 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
302 area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
303 normal = normal * area !Scale normal by area
304 this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
305 this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
306 this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
307 end do
308
309 if (neko_bcknd_device .eq. 1 .and. &
310 (unique_point_idx%num_entries() .gt. 0 )) then
311 call device_map(this%unique_mask, this%unique_mask_d, &
312 size(this%unique_mask))
313 call device_memcpy(this%unique_mask, this%unique_mask_d, &
314 size(this%unique_mask), host_to_device, sync = .true.)
315 call device_memcpy(this%nx%x, this%nx%x_d, &
316 this%nx%size(), host_to_device, sync = .true.)
317 call device_memcpy(this%ny%x, this%ny%x_d, &
318 this%ny%size(), host_to_device, sync = .true.)
319 call device_memcpy(this%nz%x, this%nz%x_d, &
320 this%nz%size(), host_to_device, sync = .true.)
321 end if
322
323 call unique_point_idx%free()
324
325 end subroutine facet_normal_finalize
326
329 class(facet_normal_t), target, intent(inout) :: this
330 integer :: i, htable_data, idx(4), facet
331 real(kind=rp) :: area, normal(3)
332
333 if (.not. allocated(this%unique_mask)) return
334 if (this%unique_mask(0) .eq. 0) return
335
336 do i = 1, this%unique_mask(0)
337 this%nx%x(i) = 0.0_rp
338 this%ny%x(i) = 0.0_rp
339 this%nz%x(i) = 0.0_rp
340 end do
341
342 do i = 1, this%facet_node_msk(0)
343 htable_data = this%msk_to_unique(i)
344 facet = this%facet(i)
345
346 idx = nonlinear_index(this%facet_node_msk(i), this%Xh%lx, this%Xh%lx, &
347 this%Xh%lx)
348 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
349 area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
350 normal = normal * area !Scale normal by area
351 this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
352 this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
353 this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
354 end do
355
356 if (neko_bcknd_device .eq. 1) then
357 call device_memcpy(this%nx%x, this%nx%x_d, &
358 this%nx%size(), host_to_device, sync = .false.)
359 call device_memcpy(this%ny%x, this%ny%x_d, &
360 this%ny%size(), host_to_device, sync = .false.)
361 call device_memcpy(this%nz%x, this%nz%x_d, &
362 this%nz%size(), host_to_device, sync = .true.)
363 end if
364
365 end subroutine facet_normal_recompute_normals
366
367end module facet_normal
__inline__ __device__ void nonlinear_index(const int idx, const int lx, int *index)
Definition bc_utils.h:44
Map a Fortran array to a device (allocate and associate)
Definition device.F90:83
Copy data between host and device (or device and device)
Definition device.F90:72
Unmap a Fortran array from a device (deassociate and free)
Definition device.F90:89
Defines a boundary condition.
Definition bc.f90:34
integer, parameter, public bc_dirichlet
Supported boundary condition types. The values are set in order of precedence for global resolution....
Definition bc.f90:66
Coefficients.
Definition coef.f90:34
subroutine, public device_masked_scatter_copy_0(a_d, b_d, mask_d, n, n_mask, strm)
Scatter a masked vector .
subroutine, public device_col2(a_d, b_d, n, strm)
Vector multiplication .
subroutine, public device_masked_gather_copy_0(a_d, b_d, mask_d, n, n_mask, strm)
Gather a masked vector .
Device abstraction, common interface for various accelerators.
Definition device.F90:34
integer, parameter, public host_to_device
Definition device.F90:48
integer, parameter, public device_to_host
Definition device.F90:48
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
Definition device.F90:52
Dirichlet condition applied in the facet normal direction.
subroutine facet_normal_init_from_components(this, coef)
Constructor from components.
subroutine facet_normal_apply_vector_dev(this, x_d, y_d, z_d, time, strong, strm)
No-op vector apply on device.
subroutine facet_normal_apply_scalar_dev(this, x_d, time, strong, strm)
No-op scalar apply on device.
subroutine facet_normal_init(this, coef, json)
Constructor.
subroutine facet_normal_recompute_normals(this)
Recompute area-weighted normals from the current mesh.
subroutine facet_normal_apply_scalar(this, x, n, time, strong)
No-op scalar apply.
subroutine facet_normal_finalize(this)
Finalize.
subroutine facet_normal_apply_surfvec_dev(this, x_d, y_d, z_d, u_d, v_d, w_d, time, strm)
Apply in facet normal direction (vector valued, device version)
subroutine facet_normal_apply_surfvec(this, x, y, z, u, v, w, n, time)
Apply in facet normal direction (vector valued)
subroutine facet_normal_apply_vector(this, x, y, z, n, time, strong)
No-op vector apply.
subroutine facet_normal_free(this)
Destructor.
Implements a hash table ADT.
Definition htable.f90:52
Definition math.f90:60
subroutine, public cfill_mask(a, c, n, mask, n_mask)
Fill a constant to a masked vector. .
Definition math.f90:491
Build configurations.
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Module with things related to the simulation time.
Utilities.
Definition utils.f90:35
Defines a vector.
Definition vector.f90:34
Base type for a boundary condition.
Definition bc.f90:72
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:93
Dirichlet condition in facet normal direction.
Integer based hash table.
Definition htable.f90:102
A struct that contains all info about the time, expand as needed.