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 procedure, pass(this) :: apply_surfvec_sub => &
69 ! > Recompute normals
70 procedure, pass(this) :: recompute_normals => facet_normal_recompute_normals
72 procedure, pass(this) :: init => facet_normal_init
74 procedure, pass(this) :: init_from_components => &
77 procedure, pass(this) :: free => facet_normal_free
79 procedure, pass(this) :: finalize => facet_normal_finalize
80 end type facet_normal_t
81
82contains
83
87 subroutine facet_normal_init(this, coef, json)
88 class(facet_normal_t), intent(inout), target :: this
89 type(coef_t), target, intent(in) :: coef
90 type(json_file), intent(inout) :: json
91
92 call this%init_from_components(coef)
93 end subroutine facet_normal_init
94
97 subroutine facet_normal_init_from_components(this, coef)
98 class(facet_normal_t), intent(inout), target :: this
99 type(coef_t), target, intent(in) :: coef
100
101 call this%init_base(coef)
102 this%bc_type = bc_dirichlet
104
106 subroutine facet_normal_apply_scalar(this, x, n, time, strong)
107 class(facet_normal_t), intent(inout) :: this
108 integer, intent(in) :: n
109 real(kind=rp), intent(inout), dimension(n) :: x
110 type(time_state_t), intent(in), optional :: time
111 logical, intent(in), optional :: strong
112 end subroutine facet_normal_apply_scalar
113
115 subroutine facet_normal_apply_scalar_dev(this, x_d, time, strong, strm)
116 class(facet_normal_t), intent(inout), target :: this
117 type(c_ptr), intent(inout) :: x_d
118 type(time_state_t), intent(in), optional :: time
119 logical, intent(in), optional :: strong
120 type(c_ptr), intent(inout) :: strm
121
122 end subroutine facet_normal_apply_scalar_dev
123
125 subroutine facet_normal_apply_vector_dev(this, x_d, y_d, z_d, time, &
126 strong, strm)
127 class(facet_normal_t), intent(inout), target :: this
128 type(c_ptr), intent(inout) :: x_d
129 type(c_ptr), intent(inout) :: y_d
130 type(c_ptr), intent(inout) :: z_d
131 type(time_state_t), intent(in), optional :: time
132 logical, intent(in), optional :: strong
133 type(c_ptr), intent(inout) :: strm
134
135 end subroutine facet_normal_apply_vector_dev
136
138 subroutine facet_normal_apply_vector(this, x, y, z, n, time, strong)
139 class(facet_normal_t), intent(inout) :: this
140 integer, intent(in) :: n
141 real(kind=rp), intent(inout), dimension(n) :: x
142 real(kind=rp), intent(inout), dimension(n) :: y
143 real(kind=rp), intent(inout), dimension(n) :: z
144 type(time_state_t), intent(in), optional :: time
145 logical, intent(in), optional :: strong
146 end subroutine facet_normal_apply_vector
147
149 subroutine facet_normal_apply_surfvec(this, x, y, z, u, v, w, n, time)
150 class(facet_normal_t), intent(in) :: this
151 integer, intent(in) :: n
152 real(kind=rp), intent(inout), dimension(n) :: x
153 real(kind=rp), intent(inout), dimension(n) :: y
154 real(kind=rp), intent(inout), dimension(n) :: z
155 real(kind=rp), intent(inout), dimension(n) :: u
156 real(kind=rp), intent(inout), dimension(n) :: v
157 real(kind=rp), intent(inout), dimension(n) :: w
158 type(time_state_t), intent(in), optional :: time
159 integer :: i, m, k
160
161 m = this%unique_mask(0)
162 ! Since apply_surfvec is called outside of the parallel region, we
163 ! need to open a separate parallel region here
164 !$omp parallel do private(k)
165 do i = 1, m
166 k = this%unique_mask(i)
167 x(k) = u(k) * this%nx%x(i)
168 y(k) = v(k) * this%ny%x(i)
169 z(k) = w(k) * this%nz%x(i)
170 end do
171 !$omp end parallel do
172
173 end subroutine facet_normal_apply_surfvec
174
197 subroutine facet_normal_apply_surfvec_sub(this, res, u, v, w, c, n)
198 class(facet_normal_t), intent(in) :: this
199 integer, intent(in) :: n
200 real(kind=rp), intent(inout), dimension(n) :: res
201 real(kind=rp), intent(in), dimension(n) :: u
202 real(kind=rp), intent(in), dimension(n) :: v
203 real(kind=rp), intent(in), dimension(n) :: w
204 real(kind=rp), intent(in) :: c
205 integer :: i, m, k
206
207 m = this%unique_mask(0)
208 !$omp do
209 do i = 1, m
210 k = this%unique_mask(i)
211 res(k) = res(k) - c * (u(k) * this%nx%x(i) &
212 + v(k) * this%ny%x(i) + w(k) * this%nz%x(i))
213 end do
214 !$omp end do
215
216 end subroutine facet_normal_apply_surfvec_sub
217
219 subroutine facet_normal_apply_surfvec_dev(this, x_d, y_d, z_d, &
220 u_d, v_d, w_d, time, strm)
221 class(facet_normal_t), intent(in), target :: this
222 type(c_ptr) :: x_d, y_d, z_d, u_d, v_d, w_d
223 type(time_state_t), intent(in), optional :: time
224 type(c_ptr), optional :: strm
225 type(c_ptr) :: strm_
226 integer :: n, m
227
228 n = this%coef%dof%size()
229 m = this%unique_mask(0)
230
231 if (present(strm)) then
232 strm_ = strm
233 else
234 strm_ = glb_cmd_queue
235 end if
236
237 if (m .gt. 0) then
238 call device_masked_gather_copy_0(this%work%x_d, u_d, &
239 this%unique_mask_d, n, m, strm_)
240 call device_col2(this%work%x_d, this%nx%x_d, m, strm_)
241 call device_masked_scatter_copy_0(x_d, this%work%x_d, &
242 this%unique_mask_d, n, m, strm_)
243 call device_masked_gather_copy_0(this%work%x_d, v_d, &
244 this%unique_mask_d, n, m, strm_)
245 call device_col2(this%work%x_d, this%ny%x_d, m, strm_)
246 call device_masked_scatter_copy_0(y_d, this%work%x_d, &
247 this%unique_mask_d, n, m, strm_)
248 call device_masked_gather_copy_0(this%work%x_d, w_d, &
249 this%unique_mask_d, n, m, strm_)
250 call device_col2(this%work%x_d, this%nz%x_d, m, strm)
251 call device_masked_scatter_copy_0(z_d, this%work%x_d, &
252 this%unique_mask_d, n, m, strm_)
253 end if
254
255 end subroutine facet_normal_apply_surfvec_dev
256
258 subroutine facet_normal_free(this)
259 class(facet_normal_t), target, intent(inout) :: this
260
261 call this%free_base()
262 if (allocated(this%unique_mask)) then
263 if (neko_bcknd_device .eq. 1) then
264 call device_unmap(this%unique_mask, this%unique_mask_d)
265 end if
266 deallocate(this%unique_mask)
267 end if
268 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
269
270 call this%nx%free()
271 call this%ny%free()
272 call this%nz%free()
273 call this%work%free()
274
275 end subroutine facet_normal_free
276
278 subroutine facet_normal_finalize(this)
279 class(facet_normal_t), target, intent(inout) :: this
280 type(htable_i4_t) :: unique_point_idx
281 integer :: htable_data, rcode, i, j, idx(4), facet
282 real(kind=rp) :: area, normal(3)
283
284 ! Here and in recompute_normals(), which only runs after this
285 call this%coef%require_facets('facet_normal')
286
287 call this%finalize_base()
288
289 ! This part is purely needed to ensure that contributions
290 ! for all faces a point is on is properly summed up.
291 ! If one simply uses the original mask, if a point is on a corner
292 ! where both faces are on the boundary
293 ! one will only get the contribution from one face, not both
294 ! We solve this by adding up the normals of both faces for these points
295 ! and storing this sum in this%nx, this%ny, this%nz.
296 ! As both contributions are added already,
297 ! we also ensure that we only visit each point once
298 ! and create a new mask with only unique points (this%unique_mask).
299 if (allocated(this%unique_mask)) then
300 if (neko_bcknd_device .eq. 1) then
301 call device_unmap(this%unique_mask, this%unique_mask_d)
302 end if
303 deallocate(this%unique_mask)
304 end if
305 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
306
307 call unique_point_idx%init(this%facet_node_msk(0), htable_data)
308 j = 0
309 do i = 1, this%facet_node_msk(0)
310 if (unique_point_idx%get(this%facet_node_msk(i), &
311 htable_data) .ne. 0) then
312 j = j + 1
313 htable_data = j
314 call unique_point_idx%set(this%facet_node_msk(i), j)
315 end if
316 end do
317
318 ! Only allocate work vectors if size is non-zero
319 if (unique_point_idx%num_entries() .gt. 0 ) then
320 call this%nx%init(unique_point_idx%num_entries())
321 call this%ny%init(unique_point_idx%num_entries())
322 call this%nz%init(unique_point_idx%num_entries())
323 call this%work%init(unique_point_idx%num_entries())
324 end if
325 allocate(this%unique_mask(0:unique_point_idx%num_entries()))
326 allocate(this%msk_to_unique(this%facet_node_msk(0)))
327
328 this%unique_mask(0) = unique_point_idx%num_entries()
329 do i = 1, this%unique_mask(0)
330 this%unique_mask(i) = 0
331 end do
332
333
334 do i = 1, this%facet_node_msk(0)
335 rcode = unique_point_idx%get(this%facet_node_msk(i), htable_data)
336 if (rcode .ne. 0) call neko_error("Facet normal: htable get failed.")
337 this%unique_mask(htable_data) = this%facet_node_msk(i)
338
339 ! Save the slot so recompute_normals can use it without the hash table.
340 this%msk_to_unique(i) = htable_data
341 facet = this%facet(i)
342
343 idx = nonlinear_index(this%facet_node_msk(i), this%Xh%lx, this%Xh%lx, &
344 this%Xh%lx)
345 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
346 area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
347 normal = normal * area !Scale normal by area
348 this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
349 this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
350 this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
351 end do
352
353 if (neko_bcknd_device .eq. 1 .and. &
354 (unique_point_idx%num_entries() .gt. 0 )) then
355 call device_map(this%unique_mask, this%unique_mask_d, &
356 size(this%unique_mask))
357 call device_memcpy(this%unique_mask, this%unique_mask_d, &
358 size(this%unique_mask), host_to_device, sync = .true.)
359 call device_memcpy(this%nx%x, this%nx%x_d, &
360 this%nx%size(), host_to_device, sync = .true.)
361 call device_memcpy(this%ny%x, this%ny%x_d, &
362 this%ny%size(), host_to_device, sync = .true.)
363 call device_memcpy(this%nz%x, this%nz%x_d, &
364 this%nz%size(), host_to_device, sync = .true.)
365 end if
366
367 call unique_point_idx%free()
368
369 end subroutine facet_normal_finalize
370
373 class(facet_normal_t), target, intent(inout) :: this
374 integer :: i, htable_data, idx(4), facet
375 real(kind=rp) :: area, normal(3)
376
377 if (.not. allocated(this%unique_mask)) return
378 if (this%unique_mask(0) .eq. 0) return
379
380 do i = 1, this%unique_mask(0)
381 this%nx%x(i) = 0.0_rp
382 this%ny%x(i) = 0.0_rp
383 this%nz%x(i) = 0.0_rp
384 end do
385
386 do i = 1, this%facet_node_msk(0)
387 htable_data = this%msk_to_unique(i)
388 facet = this%facet(i)
389
390 idx = nonlinear_index(this%facet_node_msk(i), this%Xh%lx, this%Xh%lx, &
391 this%Xh%lx)
392 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
393 area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
394 normal = normal * area !Scale normal by area
395 this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
396 this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
397 this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
398 end do
399
400 if (neko_bcknd_device .eq. 1) then
401 call device_memcpy(this%nx%x, this%nx%x_d, &
402 this%nx%size(), host_to_device, sync = .false.)
403 call device_memcpy(this%ny%x, this%ny%x_d, &
404 this%ny%size(), host_to_device, sync = .false.)
405 call device_memcpy(this%nz%x, this%nz%x_d, &
406 this%nz%size(), host_to_device, sync = .true.)
407 end if
408
409 end subroutine facet_normal_recompute_normals
410
411end 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:67
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_surfvec_sub(this, res, u, v, w, c, n)
Subtract the normal projection of a vector from a scalar field, on the facet nodes only.
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:492
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:73
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:135
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.