Neko 1.99.6
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
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)
101
103 subroutine facet_normal_apply_scalar(this, x, n, time, strong)
104 class(facet_normal_t), intent(inout) :: this
105 integer, intent(in) :: n
106 real(kind=rp), intent(inout), dimension(n) :: x
107 type(time_state_t), intent(in), optional :: time
108 logical, intent(in), optional :: strong
109 end subroutine facet_normal_apply_scalar
110
112 subroutine facet_normal_apply_scalar_dev(this, x_d, time, strong, strm)
113 class(facet_normal_t), intent(inout), target :: this
114 type(c_ptr), intent(inout) :: x_d
115 type(time_state_t), intent(in), optional :: time
116 logical, intent(in), optional :: strong
117 type(c_ptr), intent(inout) :: strm
118
119 end subroutine facet_normal_apply_scalar_dev
120
122 subroutine facet_normal_apply_vector_dev(this, x_d, y_d, z_d, time, &
123 strong, strm)
124 class(facet_normal_t), intent(inout), target :: this
125 type(c_ptr), intent(inout) :: x_d
126 type(c_ptr), intent(inout) :: y_d
127 type(c_ptr), intent(inout) :: z_d
128 type(time_state_t), intent(in), optional :: time
129 logical, intent(in), optional :: strong
130 type(c_ptr), intent(inout) :: strm
131
132 end subroutine facet_normal_apply_vector_dev
133
135 subroutine facet_normal_apply_vector(this, x, y, z, n, time, strong)
136 class(facet_normal_t), intent(inout) :: this
137 integer, intent(in) :: n
138 real(kind=rp), intent(inout), dimension(n) :: x
139 real(kind=rp), intent(inout), dimension(n) :: y
140 real(kind=rp), intent(inout), dimension(n) :: z
141 type(time_state_t), intent(in), optional :: time
142 logical, intent(in), optional :: strong
143 end subroutine facet_normal_apply_vector
144
146 subroutine facet_normal_apply_surfvec(this, x, y, z, u, v, w, n, time)
147 class(facet_normal_t), intent(in) :: this
148 integer, intent(in) :: n
149 real(kind=rp), intent(inout), dimension(n) :: x
150 real(kind=rp), intent(inout), dimension(n) :: y
151 real(kind=rp), intent(inout), dimension(n) :: z
152 real(kind=rp), intent(inout), dimension(n) :: u
153 real(kind=rp), intent(inout), dimension(n) :: v
154 real(kind=rp), intent(inout), dimension(n) :: w
155 type(time_state_t), intent(in), optional :: time
156 integer :: i, m, k, idx(4), facet
157 real(kind=rp) :: normal(3), area
158
159 m = this%unique_mask(0)
160 ! Since apply_surfvec is called outside of the parallel region, we
161 ! need to open a separate parallel region here
162 !$omp parallel do
163 do i = 1, m
164 k = this%unique_mask(i)
165 x(k) = u(k) * this%nx%x(i)
166 y(k) = v(k) * this%ny%x(i)
167 z(k) = w(k) * this%nz%x(i)
168 end do
169 !$omp end parallel do
170
171 end subroutine facet_normal_apply_surfvec
172
174 subroutine facet_normal_apply_surfvec_dev(this, x_d, y_d, z_d, &
175 u_d, v_d, w_d, time, strm)
176 class(facet_normal_t), intent(in), target :: this
177 type(c_ptr) :: x_d, y_d, z_d, u_d, v_d, w_d
178 type(time_state_t), intent(in), optional :: time
179 type(c_ptr), optional :: strm
180 type(c_ptr) :: strm_
181 integer :: n, m
182
183 n = this%coef%dof%size()
184 m = this%unique_mask(0)
185
186 if (present(strm)) then
187 strm_ = strm
188 else
189 strm_ = glb_cmd_queue
190 end if
191
192 if (m .gt. 0) then
193 call device_masked_gather_copy_0(this%work%x_d, u_d, &
194 this%unique_mask_d, n, m, strm_)
195 call device_col2(this%work%x_d, this%nx%x_d, m, strm_)
196 call device_masked_scatter_copy_0(x_d, this%work%x_d, &
197 this%unique_mask_d, n, m, strm_)
198 call device_masked_gather_copy_0(this%work%x_d, v_d, &
199 this%unique_mask_d, n, m, strm_)
200 call device_col2(this%work%x_d, this%ny%x_d, m, strm_)
201 call device_masked_scatter_copy_0(y_d, this%work%x_d, &
202 this%unique_mask_d, n, m, strm_)
203 call device_masked_gather_copy_0(this%work%x_d, w_d, &
204 this%unique_mask_d, n, m, strm_)
205 call device_col2(this%work%x_d, this%nz%x_d, m, strm)
206 call device_masked_scatter_copy_0(z_d, this%work%x_d, &
207 this%unique_mask_d, n, m, strm_)
208 end if
209
210 end subroutine facet_normal_apply_surfvec_dev
211
213 subroutine facet_normal_free(this)
214 class(facet_normal_t), target, intent(inout) :: this
215
216 call this%free_base()
217 if (allocated(this%unique_mask)) then
218 if (neko_bcknd_device .eq. 1) then
219 call device_unmap(this%unique_mask, this%unique_mask_d)
220 end if
221 deallocate(this%unique_mask)
222 end if
223 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
224
225 call this%nx%free()
226 call this%ny%free()
227 call this%nz%free()
228 call this%work%free()
229
230 end subroutine facet_normal_free
231
233 subroutine facet_normal_finalize(this, only_facets)
234 class(facet_normal_t), target, intent(inout) :: this
235 logical, optional, intent(in) :: only_facets
236 logical :: only_facets_
237 type(htable_i4_t) :: unique_point_idx
238 integer :: htable_data, rcode, i, j, idx(4), facet
239 real(kind=rp) :: area, normal(3)
240
241 if (present(only_facets)) then
242 if (.not. only_facets) then
243 call neko_error("For facet_normal_t, only_facets has to be true.")
244 end if
245 end if
246
247 call this%finalize_base(.true.)
248 ! This part is purely needed to ensure that contributions
249 ! for all faces a point is on is properly summed up.
250 ! If one simply uses the original mask, if a point is on a corner
251 ! where both faces are on the boundary
252 ! one will only get the contribution from one face, not both
253 ! We solve this by adding up the normals of both faces for these points
254 ! and storing this sum in this%nx, this%ny, this%nz.
255 ! As both contrbutions are added already,
256 ! we also ensure that we only visit each point once
257 ! and create a new mask with only unique points (this%unique_mask).
258 if (allocated(this%unique_mask)) then
259 if (neko_bcknd_device .eq. 1) then
260 call device_unmap(this%unique_mask, this%unique_mask_d)
261 end if
262 deallocate(this%unique_mask)
263 end if
264 if (allocated(this%msk_to_unique)) deallocate(this%msk_to_unique)
265
266 call unique_point_idx%init(this%msk(0), htable_data)
267 j = 0
268 do i = 1, this%msk(0)
269 if (unique_point_idx%get(this%msk(i), htable_data) .ne. 0) then
270 j = j + 1
271 htable_data = j
272 call unique_point_idx%set(this%msk(i), j)
273 end if
274 end do
275
276 ! Only allocate work vectors if size is non-zero
277 if (unique_point_idx%num_entries() .gt. 0 ) then
278 call this%nx%init(unique_point_idx%num_entries())
279 call this%ny%init(unique_point_idx%num_entries())
280 call this%nz%init(unique_point_idx%num_entries())
281 call this%work%init(unique_point_idx%num_entries())
282 end if
283 allocate(this%unique_mask(0:unique_point_idx%num_entries()))
284 allocate(this%msk_to_unique(this%msk(0)))
285
286 this%unique_mask(0) = unique_point_idx%num_entries()
287 do i = 1, this%unique_mask(0)
288 this%unique_mask(i) = 0
289 end do
290
291
292 do i = 1, this%msk(0)
293 rcode = unique_point_idx%get(this%msk(i), htable_data)
294 if (rcode .ne. 0) call neko_error("Facet normal: htable get failed.")
295 this%unique_mask(htable_data) = this%msk(i)
296
297 ! Save the slot so recompute_normals can use it without the hash table.
298 this%msk_to_unique(i) = htable_data
299 facet = this%facet(i)
300
301 idx = nonlinear_index(this%msk(i), this%Xh%lx, this%Xh%lx, this%Xh%lx)
302 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
303 area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
304 normal = normal * area !Scale normal by area
305 this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
306 this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
307 this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
308 end do
309
310 if (neko_bcknd_device .eq. 1 .and. &
311 (unique_point_idx%num_entries() .gt. 0 )) then
312 call device_map(this%unique_mask, this%unique_mask_d, &
313 size(this%unique_mask))
314 call device_memcpy(this%unique_mask, this%unique_mask_d, &
315 size(this%unique_mask), host_to_device, sync = .true.)
316 call device_memcpy(this%nx%x, this%nx%x_d, &
317 this%nx%size(), host_to_device, sync = .true.)
318 call device_memcpy(this%ny%x, this%ny%x_d, &
319 this%ny%size(), host_to_device, sync = .true.)
320 call device_memcpy(this%nz%x, this%nz%x_d, &
321 this%nz%size(), host_to_device, sync = .true.)
322 end if
323
324 call unique_point_idx%free()
325
326 end subroutine facet_normal_finalize
327
330 class(facet_normal_t), target, intent(inout) :: this
331 integer :: i, htable_data, idx(4), facet
332 real(kind=rp) :: area, normal(3)
333
334 if (.not. allocated(this%unique_mask)) return
335 if (this%unique_mask(0) .eq. 0) return
336
337 do i = 1, this%unique_mask(0)
338 this%nx%x(i) = 0.0_rp
339 this%ny%x(i) = 0.0_rp
340 this%nz%x(i) = 0.0_rp
341 end do
342
343 do i = 1, this%msk(0)
344 htable_data = this%msk_to_unique(i)
345 facet = this%facet(i)
346
347 idx = nonlinear_index(this%msk(i), this%Xh%lx, this%Xh%lx, 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
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_finalize(this, only_facets)
Finalize.
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_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:488
Build configurations.
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:12
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:62
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:63
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.