44 use json_module,
only : json_file
45 use,
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
55 integer,
allocatable :: unique_mask(:)
56 integer,
allocatable :: msk_to_unique(:)
57 type(c_ptr) :: unique_mask_d = c_null_ptr
65 procedure, pass(this) :: apply_surfvec_dev => &
67 procedure, pass(this) :: apply_surfvec_sub => &
74 procedure, pass(this) :: init_from_components => &
89 type(
coef_t),
target,
intent(in) :: coef
90 type(json_file),
intent(inout) :: json
92 call this%init_from_components(coef)
99 type(
coef_t),
target,
intent(in) :: coef
101 call this%init_base(coef)
108 integer,
intent(in) :: n
109 real(kind=
rp),
intent(inout),
dimension(n) :: x
111 logical,
intent(in),
optional :: strong
117 type(c_ptr),
intent(inout) :: x_d
119 logical,
intent(in),
optional :: strong
120 type(c_ptr),
intent(inout) :: strm
128 type(c_ptr),
intent(inout) :: x_d
129 type(c_ptr),
intent(inout) :: y_d
130 type(c_ptr),
intent(inout) :: z_d
132 logical,
intent(in),
optional :: strong
133 type(c_ptr),
intent(inout) :: strm
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
145 logical,
intent(in),
optional :: strong
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
161 m = this%unique_mask(0)
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)
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
207 m = this%unique_mask(0)
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))
220 u_d, v_d, w_d, time, strm)
222 type(c_ptr) :: x_d, y_d, z_d, u_d, v_d, w_d
224 type(c_ptr),
optional :: strm
228 n = this%coef%dof%size()
229 m = this%unique_mask(0)
231 if (
present(strm))
then
239 this%unique_mask_d, n, m, strm_)
240 call device_col2(this%work%x_d, this%nx%x_d, m, strm_)
242 this%unique_mask_d, n, m, strm_)
244 this%unique_mask_d, n, m, strm_)
245 call device_col2(this%work%x_d, this%ny%x_d, m, strm_)
247 this%unique_mask_d, n, m, strm_)
249 this%unique_mask_d, n, m, strm_)
250 call device_col2(this%work%x_d, this%nz%x_d, m, strm)
252 this%unique_mask_d, n, m, strm_)
261 call this%free_base()
262 if (
allocated(this%unique_mask))
then
266 deallocate(this%unique_mask)
268 if (
allocated(this%msk_to_unique))
deallocate(this%msk_to_unique)
273 call this%work%free()
281 integer :: htable_data, rcode, i, j, idx(4), facet
282 real(kind=
rp) :: area, normal(3)
285 call this%coef%require_facets(
'facet_normal')
287 call this%finalize_base()
299 if (
allocated(this%unique_mask))
then
303 deallocate(this%unique_mask)
305 if (
allocated(this%msk_to_unique))
deallocate(this%msk_to_unique)
307 call unique_point_idx%init(this%facet_node_msk(0), htable_data)
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
314 call unique_point_idx%set(this%facet_node_msk(i), j)
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())
325 allocate(this%unique_mask(0:unique_point_idx%num_entries()))
326 allocate(this%msk_to_unique(this%facet_node_msk(0)))
328 this%unique_mask(0) = unique_point_idx%num_entries()
329 do i = 1, this%unique_mask(0)
330 this%unique_mask(i) = 0
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)
340 this%msk_to_unique(i) = htable_data
341 facet = this%facet(i)
343 idx =
nonlinear_index(this%facet_node_msk(i), this%Xh%lx, 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
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)
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))
367 call unique_point_idx%free()
374 integer :: i, htable_data, idx(4), facet
375 real(kind=
rp) :: area, normal(3)
377 if (.not.
allocated(this%unique_mask))
return
378 if (this%unique_mask(0) .eq. 0)
return
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
386 do i = 1, this%facet_node_msk(0)
387 htable_data = this%msk_to_unique(i)
388 facet = this%facet(i)
390 idx =
nonlinear_index(this%facet_node_msk(i), this%Xh%lx, 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
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)
__inline__ __device__ void nonlinear_index(const int idx, const int lx, int *index)
Map a Fortran array to a device (allocate and associate)
Copy data between host and device (or device and device)
Unmap a Fortran array from a device (deassociate and free)
Defines a boundary condition.
integer, parameter, public bc_dirichlet
Supported boundary condition types. The values are set in order of precedence for global resolution....
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.
integer, parameter, public host_to_device
integer, parameter, public device_to_host
type(c_ptr), bind(C), public glb_cmd_queue
Global command queue.
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.
subroutine, public cfill_mask(a, c, n, mask, n_mask)
Fill a constant to a masked vector. .
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Module with things related to the simulation time.
Base type for a boundary condition.
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Dirichlet condition in facet normal direction.
Integer based hash table.
A struct that contains all info about the time, expand as needed.