45  use json_module, 
only : json_file
 
   46  use, 
intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
 
   56     integer, 
allocatable :: unique_mask(:)
 
   57     type(c_ptr) :: unique_mask_d = c_null_ptr
 
   65     procedure, pass(this) :: apply_surfvec_dev => &
 
   70     procedure, pass(this) :: init_from_components => &
 
 
   85    type(
coef_t), 
target, 
intent(in) :: coef
 
   86    type(json_file), 
intent(inout) ::json
 
   88    call this%init_from_components(coef)
 
 
   95    type(
coef_t), 
target, 
intent(in) :: coef
 
   97    call this%init_base(coef)
 
 
  103    integer, 
intent(in) :: n
 
  104    real(kind=
rp), 
intent(inout), 
dimension(n) :: x
 
  106    logical, 
intent(in), 
optional :: strong
 
 
  112    type(c_ptr),
intent(inout) :: x_d
 
  114    logical, 
intent(in), 
optional :: strong
 
  115    type(c_ptr), 
intent(inout) :: strm
 
 
  123    type(c_ptr), 
intent(inout) :: x_d
 
  124    type(c_ptr), 
intent(inout) :: y_d
 
  125    type(c_ptr), 
intent(inout) :: z_d
 
  127    logical, 
intent(in), 
optional :: strong
 
  128    type(c_ptr), 
intent(inout) :: strm
 
 
  135    integer, 
intent(in) :: n
 
  136    real(kind=
rp), 
intent(inout), 
dimension(n) :: x
 
  137    real(kind=
rp), 
intent(inout), 
dimension(n) :: y
 
  138    real(kind=
rp), 
intent(inout), 
dimension(n) :: z
 
  140    logical, 
intent(in), 
optional :: strong
 
 
  146    integer, 
intent(in) :: n
 
  147    real(kind=
rp), 
intent(inout), 
dimension(n) :: x
 
  148    real(kind=
rp), 
intent(inout), 
dimension(n) :: y
 
  149    real(kind=
rp), 
intent(inout), 
dimension(n) :: z
 
  150    real(kind=
rp), 
intent(inout), 
dimension(n) :: u
 
  151    real(kind=
rp), 
intent(inout), 
dimension(n) :: v
 
  152    real(kind=
rp), 
intent(inout), 
dimension(n) :: w
 
  154    integer :: i, m, k, idx(4), facet
 
  155    real(kind=
rp) :: normal(3), area
 
  157    m = this%unique_mask(0)
 
  160       k = this%unique_mask(i)
 
  161       x(k) = u(k) * this%nx%x(i)
 
  162       y(k) = v(k) * this%ny%x(i)
 
  163       z(k) = w(k) * this%nz%x(i)
 
 
  170       u_d, v_d, w_d, time, strm)
 
  172    type(c_ptr) :: x_d, y_d, z_d, u_d, v_d, w_d
 
  174    type(c_ptr), 
optional :: strm
 
  178    n = this%coef%dof%size()
 
  179    m = this%unique_mask(0)
 
  181    if (
present(strm)) 
then 
  190       call device_col2(this%work%x_d, this%nx%x_d, m, strm_)
 
  192            this%unique_mask_d, n, m, strm_)
 
  195       call device_col2(this%work%x_d, this%ny%x_d, m, strm_)
 
  197            this%unique_mask_d, n, m, strm_)
 
  200       call device_col2(this%work%x_d, this%nz%x_d, m, strm)
 
  202            this%unique_mask_d, n, m, strm_)
 
 
  211    call this%free_base()
 
  212    if (
allocated(this%unique_mask)) 
then 
  213       deallocate(this%unique_mask)
 
  215    if (c_associated(this%unique_mask_d)) 
then 
  222    call this%work%free()
 
 
  229    logical, 
optional, 
intent(in) :: only_facets
 
  230    logical :: only_facets_
 
  232    integer :: htable_data, rcode, i, j, idx(4), facet
 
  233    real(kind=
rp) :: area, normal(3)
 
  235    if (
present(only_facets)) 
then 
  236       if (only_facets .eqv. .false.) 
then 
  237          call neko_error(
"For facet_normal_t, only_facets has to be true.")
 
  241    call this%finalize_base(.true.)
 
  252    if (
allocated(this%unique_mask)) 
then 
  253       deallocate(this%unique_mask)
 
  255    if (c_associated(this%unique_mask_d)) 
then 
  259    call unique_point_idx%init(this%msk(0), htable_data)
 
  261    do i = 1, this%msk(0)
 
  262       if (unique_point_idx%get(this%msk(i),htable_data) .ne. 0) 
then 
  265          call unique_point_idx%set(this%msk(i), j)
 
  270    if (unique_point_idx%num_entries() .gt. 0 ) 
then 
  271       call this%nx%init(unique_point_idx%num_entries())
 
  272       call this%ny%init(unique_point_idx%num_entries())
 
  273       call this%nz%init(unique_point_idx%num_entries())
 
  274       call this%work%init(unique_point_idx%num_entries())
 
  276    allocate(this%unique_mask(0:unique_point_idx%num_entries()))
 
  278    this%unique_mask(0) = unique_point_idx%num_entries()
 
  279    do i = 1, this%unique_mask(0)
 
  280       this%unique_mask(i) = 0
 
  284    do i = 1, this%msk(0)
 
  285       rcode = unique_point_idx%get(this%msk(i), htable_data)
 
  286       if (rcode .ne. 0) 
call neko_error(
"Facet normal: htable get failed.")
 
  287       this%unique_mask(htable_data) = this%msk(i)
 
  288       facet = this%facet(i)
 
  291       normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), facet)
 
  292       area = this%coef%get_area(idx(1), idx(2), idx(3), idx(4), facet)
 
  293       normal = normal * area 
 
  294       this%nx%x(htable_data) = this%nx%x(htable_data) + normal(1)
 
  295       this%ny%x(htable_data) = this%ny%x(htable_data) + normal(2)
 
  296       this%nz%x(htable_data) = this%nz%x(htable_data) + normal(3)
 
  300         (unique_point_idx%num_entries() .gt. 0 )) 
then 
  301       call device_map(this%unique_mask, this%unique_mask_d, &
 
  302            size(this%unique_mask))
 
  313    call unique_point_idx%free()
 
 
__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)
 
Defines a boundary condition.
 
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
 
subroutine, public device_free(x_d)
Deallocate memory on the 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_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_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.
 
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.