Neko 1.99.9
A portable framework for high-order spectral element flow simulations
Loading...
Searching...
No Matches
vector_bc_projector.f90
Go to the documentation of this file.
1! Copyright (c) 2026, 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!
38 use bc, only : bc_t, bc_dirichlet
39 use bc_list, only : bc_list_t
40 use mixed_bc, only : mixed_bc_t
41 use mask, only : mask_t
42 use coefs, only : coef_t
43 use dofmap, only : dofmap_t
44 use field, only : field_t
45 use field_list, only : field_list_t
46 use fld_file, only : fld_file_t
48 use htable, only : htable_i4_t
49 use logger, only : log_size
50 use matrix, only : matrix_t
53 use gs_ops, only : gs_op_add, gs_op_min
55 use num_types, only : rp
56 use tuple, only : tuple_i4_t
64 use operators, only : rotate_cyc
65 use, intrinsic :: iso_c_binding, only : c_ptr, c_null_ptr, c_associated
66 implicit none
67 private
68
70
72 type, public, abstract :: vector_bc_projector_t
73 contains
75 procedure(vector_bc_projector_init_intrf), pass(this), deferred :: init
77 procedure(vector_bc_projector_free_intrf), pass(this), deferred :: free
79 procedure(vector_bc_projector_finalize_intrf), pass(this), deferred :: &
80 finalize
82 procedure(vector_bc_projector_apply_intrf), pass(this), deferred :: apply
84 procedure(vector_bc_projector_mark_bc_intrf), pass(this), deferred :: &
85 mark_bc
87 procedure, pass(this) :: mark_bc_component => &
90 procedure(vector_bc_projector_mark_bc_list_intrf), pass(this), deferred :: &
91 mark_bc_list
93 procedure, pass(this) :: mark_bc_list_component => &
95 generic :: mark => mark_bc, mark_bc_component, mark_bc_list, &
96 mark_bc_list_component
98
112 contains
114 procedure, pass(this) :: init => segregated_vector_bc_projector_init
116 procedure, pass(this) :: free => segregated_vector_bc_projector_free
118 procedure, pass(this) :: finalize => segregated_vector_bc_projector_finalize
120 procedure, pass(this) :: mark_bc => segregated_vector_bc_projector_mark_bc
122 procedure, pass(this) :: mark_bc_component => &
125 procedure, pass(this) :: mark_bc_list => &
128 procedure, pass(this) :: mark_bc_list_component => &
131 procedure, pass(this) :: apply => segregated_vector_bc_projector_apply
133
151 type(mask_t) :: dirichlet_dof_mask
153 type(mask_t) :: mixed_dof_mask
155 type(bc_list_t), private :: bcs
157 type(coef_t), pointer, private :: coef => null()
159 type(dofmap_t), pointer, private :: dof => null()
161 real(kind=rp), allocatable :: node_type(:)
163 real(kind=rp), allocatable :: face_type(:,:)
165 integer, allocatable :: boundary_dof(:)
167 type(htable_i4_t) :: boundary_idx
169 integer, allocatable :: node_rst(:,:)
171 integer, allocatable :: edge_mid_rst(:,:)
173 integer, allocatable :: node_linear_idx(:)
175 integer, allocatable :: constraint_n(:)
177 integer, allocatable :: constraint_t1(:)
179 integer, allocatable :: constraint_t2(:)
181 type(matrix_t) :: n
183 type(matrix_t) :: t1
185 type(matrix_t) :: t2
186
188 type(c_ptr) :: constraint_n_d = c_null_ptr
190 type(c_ptr) :: constraint_t1_d = c_null_ptr
192 type(c_ptr) :: constraint_t2_d = c_null_ptr
193 contains
195 procedure, pass(this) :: init => coupled_vector_bc_projector_init
197 procedure, pass(this) :: free => coupled_vector_bc_projector_free
199 procedure, pass(this) :: mark_bc => coupled_vector_bc_projector_mark_bc
201 procedure, pass(this) :: mark_bc_list => &
204 procedure, pass(this) :: finalize => coupled_vector_bc_projector_finalize
206 procedure, pass(this) :: apply => coupled_vector_bc_projector_apply
208 procedure, pass(this) :: debug_output => &
211 procedure, pass(this) :: debug_output_normal_component => &
214 procedure, pass(this), private :: clear_masks => &
217 procedure, pass(this), private :: rebuild_masks => &
220 procedure, pass(this), private :: clear_basis => &
223 procedure, pass(this), private :: rebuild_basis => &
226
227 abstract interface
228
230 import :: vector_bc_projector_t
231 class(vector_bc_projector_t), intent(inout) :: this
232 end subroutine vector_bc_projector_free_intrf
233 end interface
234
235 abstract interface
236
238 subroutine vector_bc_projector_init_intrf(this, coef)
240 class(vector_bc_projector_t), intent(inout) :: this
241 type(coef_t), target, intent(in) :: coef
242 end subroutine vector_bc_projector_init_intrf
243 end interface
244
245 abstract interface
246
247 subroutine vector_bc_projector_finalize_intrf(this, rebuild_mask)
248 import :: vector_bc_projector_t
249 class(vector_bc_projector_t), intent(inout) :: this
250 logical, intent(in) :: rebuild_mask
252 end interface
253
254 abstract interface
255
261 subroutine vector_bc_projector_apply_intrf(this, x, y, z, n, strm)
262 import :: vector_bc_projector_t, rp, c_ptr
263 class(vector_bc_projector_t), intent(in) :: this
264 integer, intent(in) :: n
265 real(kind=rp), intent(inout) :: x(n)
266 real(kind=rp), intent(inout) :: y(n)
267 real(kind=rp), intent(inout) :: z(n)
268 type(c_ptr), intent(inout), optional :: strm
270 end interface
271
272 abstract interface
273
276 import :: vector_bc_projector_t, bc_t
277 class(vector_bc_projector_t), intent(inout) :: this
278 class(bc_t), intent(inout), target :: bc
280 end interface
281
282 abstract interface
283
287 class(vector_bc_projector_t), intent(inout) :: this
288 type(bc_list_t), intent(in) :: bclst
290 end interface
291
292contains
293
296 component)
297 class(vector_bc_projector_t), intent(inout) :: this
298 class(bc_t), intent(inout), target :: bc
299 character(len=1), intent(in) :: component
300
301 call neko_error("Component-specific marking is only supported by " // &
302 "segregated vector BC projectors.")
304
307 bclst, component)
308 class(vector_bc_projector_t), intent(inout) :: this
309 type(bc_list_t), intent(in) :: bclst
310 character(len=1), intent(in) :: component
311
312 call neko_error("Component-specific marking is only supported by " // &
313 "segregated vector BC projectors.")
315
318 class(segregated_vector_bc_projector_t), intent(inout) :: this
319 call this%x%free()
320 call this%y%free()
321 call this%z%free()
323
327 class(segregated_vector_bc_projector_t), intent(inout) :: this
328 type(coef_t), target, intent(in) :: coef
329
330 ! Just ensure we get a fresh object
331 call this%free()
333
335 subroutine segregated_vector_bc_projector_finalize(this, rebuild_mask)
336 class(segregated_vector_bc_projector_t), intent(inout) :: this
337 logical, intent(in) :: rebuild_mask
339
345 class(segregated_vector_bc_projector_t), intent(inout) :: this
346 class(bc_t), intent(inout), target :: bc
347
348 if (bc%bc_type .ne. bc_dirichlet) then
349 call neko_error("Segregated vector BC projector only accepts " // &
350 "Dirichlet boundary conditions.")
351 end if
352
353 call this%x%mark_bc(bc)
354 call this%y%mark_bc(bc)
355 call this%z%mark_bc(bc)
357
362 component)
363 class(segregated_vector_bc_projector_t), intent(inout) :: this
364 class(bc_t), intent(inout), target :: bc
365 character(len=1), intent(in) :: component
366
367 if (bc%bc_type .ne. bc_dirichlet) then
368 call neko_error("Segregated vector BC projector only accepts " // &
369 "Dirichlet boundary conditions.")
370 end if
371
372 select case (component)
373 case ('x')
374 call this%x%mark_bc(bc)
375 case ('y')
376 call this%y%mark_bc(bc)
377 case ('z')
378 call this%z%mark_bc(bc)
379 case default
380 call neko_error("Invalid component for segregated vector BC " // &
381 "projector mark.")
382 end select
384
388 class(segregated_vector_bc_projector_t), intent(inout) :: this
389 type(bc_list_t), intent(in) :: bclst
390 class(bc_t), pointer :: bc_i
391 integer :: i
392
393 do i = 1, bclst%size()
394 bc_i => bclst%get(i)
395 call this%mark(bc_i)
396 end do
398
403 bclst, component)
404 class(segregated_vector_bc_projector_t), intent(inout) :: this
405 type(bc_list_t), intent(in) :: bclst
406 character(len=1), intent(in) :: component
407 class(bc_t), pointer :: bc_i
408 integer :: i
409
410 do i = 1, bclst%size()
411 bc_i => bclst%get(i)
412 call this%mark(bc_i, component)
413 end do
415
422 subroutine segregated_vector_bc_projector_apply(this, x, y, z, n, strm)
423 class(segregated_vector_bc_projector_t), intent(in) :: this
424 integer, intent(in) :: n
425 real(kind=rp), intent(inout) :: x(n)
426 real(kind=rp), intent(inout) :: y(n)
427 real(kind=rp), intent(inout) :: z(n)
428 type(c_ptr), intent(inout), optional :: strm
429
430 call this%x%apply(x, n, strm = strm)
431 call this%y%apply(y, n, strm = strm)
432 call this%z%apply(z, n, strm = strm)
434
439 subroutine vector_bc_projector_components(this, x, y, z)
440 class(vector_bc_projector_t), target, intent(inout) :: this
441 type(scalar_bc_projector_t), pointer, intent(inout) :: x
442 type(scalar_bc_projector_t), pointer, intent(inout) :: y
443 type(scalar_bc_projector_t), pointer, intent(inout) :: z
444
445 select type (this)
447 x => this%x
448 y => this%y
449 z => this%z
450 class default
451 call neko_error("Component access is only available for " // &
452 "segregated vector BC projectors. You have likely forgotten to " // &
453 "select a coupled linear solver for velocity in the fluid " // &
454 "configuration.")
455 end select
456 end subroutine vector_bc_projector_components
457
458
459 !
460 ! Coupled projector TBPs
461 !
462
465 class(coupled_vector_bc_projector_t), intent(inout) :: this
466
467 call this%bcs%free()
468 if (allocated(this%node_rst)) deallocate(this%node_rst)
469 if (allocated(this%edge_mid_rst)) deallocate(this%edge_mid_rst)
470 if (allocated(this%node_linear_idx)) deallocate(this%node_linear_idx)
471 if (allocated(this%face_type)) deallocate(this%face_type)
472 call this%clear_masks()
473 call this%clear_basis()
474 this%constraint_n_d = c_null_ptr
475 this%constraint_t1_d = c_null_ptr
476 this%constraint_t2_d = c_null_ptr
477 nullify(this%coef)
478 nullify(this%dof)
480
486 subroutine coupled_vector_bc_projector_init(this, coef)
487 class(coupled_vector_bc_projector_t), intent(inout) :: this
488 type(coef_t), target, intent(in) :: coef
489 integer :: lx, ly, lz
490 integer :: mid_i, mid_j, mid_k
491 integer :: nface
492
493 call this%free()
494 call this%bcs%init()
495
496 this%coef => coef
497 this%dof => coef%dof
498
499 lx = coef%Xh%lx
500 ly = coef%Xh%ly
501 lz = coef%Xh%lz
502 mid_i = (lx + 1) / 2
503 mid_j = (ly + 1) / 2
504 mid_k = (lz + 1) / 2
505
506 allocate(this%node_rst(3, 8))
507 allocate(this%edge_mid_rst(3, 12))
508 allocate(this%node_linear_idx(8))
509 nface = 2 * coef%msh%gdim
510 allocate(this%face_type(nface, coef%msh%nelv))
511 this%face_type = 5.0_rp
512
513 this%node_rst(:,1) = [1, 1, 1]
514 this%node_rst(:,2) = [lx, 1, 1]
515 this%node_rst(:,3) = [1, ly, 1]
516 this%node_rst(:,4) = [lx, ly, 1]
517 this%node_rst(:,5) = [1, 1, lz]
518 this%node_rst(:,6) = [lx, 1, lz]
519 this%node_rst(:,7) = [1, ly, lz]
520 this%node_rst(:,8) = [lx, ly, lz]
521
522 this%edge_mid_rst(:,1) = [mid_i, 1, 1]
523 this%edge_mid_rst(:,2) = [mid_i, ly, 1]
524 this%edge_mid_rst(:,3) = [mid_i, 1, lz]
525 this%edge_mid_rst(:,4) = [mid_i, ly, lz]
526 this%edge_mid_rst(:,5) = [1, mid_j, 1]
527 this%edge_mid_rst(:,6) = [lx, mid_j, 1]
528 this%edge_mid_rst(:,7) = [1, mid_j, lz]
529 this%edge_mid_rst(:,8) = [lx, mid_j, lz]
530 this%edge_mid_rst(:,9) = [1, 1, mid_k]
531 this%edge_mid_rst(:,10) = [lx, 1, mid_k]
532 this%edge_mid_rst(:,11) = [1, ly, mid_k]
533 this%edge_mid_rst(:,12) = [lx, ly, mid_k]
534
535 this%node_linear_idx(1) = linear_index(1, 1, 1, 1, lx, ly, lz)
536 this%node_linear_idx(2) = linear_index(lx, 1, 1, 1, lx, ly, lz)
537 this%node_linear_idx(3) = linear_index(1, ly, 1, 1, lx, ly, lz)
538 this%node_linear_idx(4) = linear_index(lx, ly, 1, 1, lx, ly, lz)
539 this%node_linear_idx(5) = linear_index(1, 1, lz, 1, lx, ly, lz)
540 this%node_linear_idx(6) = linear_index(lx, 1, lz, 1, lx, ly, lz)
541 this%node_linear_idx(7) = linear_index(1, ly, lz, 1, lx, ly, lz)
542 this%node_linear_idx(8) = linear_index(lx, ly, lz, 1, lx, ly, lz)
544
548 class(coupled_vector_bc_projector_t), intent(inout) :: this
549 class(bc_t), intent(inout), target :: bc
550
551 if (.not. associated(this%coef)) then
552 call neko_error("Coupled vector BC projector must be initialized " // &
553 "before mark().")
554 end if
555
556 call this%bcs%append(bc)
558
562 class(coupled_vector_bc_projector_t), intent(inout) :: this
563 type(bc_list_t), intent(in) :: bclst
564 class(bc_t), pointer :: bc_i
565 integer :: i
566
567 do i = 1, bclst%size()
568 bc_i => bclst%get(i)
569 call this%mark(bc_i)
570 end do
572
578 class(coupled_vector_bc_projector_t), intent(inout) :: this
579
580 call this%dirichlet_dof_mask%free()
581 call this%mixed_dof_mask%free()
582 if (allocated(this%boundary_dof)) deallocate(this%boundary_dof)
583 if (allocated(this%node_type)) deallocate(this%node_type)
584 call this%boundary_idx%free()
585
586 if (allocated(this%constraint_n)) then
587 if (neko_bcknd_device .eq. 1 .and. &
588 c_associated(this%constraint_n_d)) then
589 call device_unmap(this%constraint_n, this%constraint_n_d)
590 end if
591 deallocate(this%constraint_n)
592 end if
593 if (allocated(this%constraint_t1)) then
594 if (neko_bcknd_device .eq. 1 .and. &
595 c_associated(this%constraint_t1_d)) then
596 call device_unmap(this%constraint_t1, this%constraint_t1_d)
597 end if
598 deallocate(this%constraint_t1)
599 end if
600 if (allocated(this%constraint_t2)) then
601 if (neko_bcknd_device .eq. 1 .and. &
602 c_associated(this%constraint_t2_d)) then
603 call device_unmap(this%constraint_t2, this%constraint_t2_d)
604 end if
605 deallocate(this%constraint_t2)
606 end if
607 this%constraint_n_d = c_null_ptr
608 this%constraint_t1_d = c_null_ptr
609 this%constraint_t2_d = c_null_ptr
611
614 class(coupled_vector_bc_projector_t), intent(inout) :: this
615
616 call this%n%free()
617 call this%t1%free()
618 call this%t2%free()
620
624 subroutine coupled_vector_bc_projector_finalize(this, rebuild_mask)
625 class(coupled_vector_bc_projector_t), intent(inout) :: this
626 logical, intent(in) :: rebuild_mask
627 if (this%bcs%size() .eq. 0) return
628
629 if (rebuild_mask) then
630 call this%rebuild_masks()
631 end if
632
633 call this%rebuild_basis()
635
643 class(coupled_vector_bc_projector_t), intent(inout) :: this
644 type(field_t), pointer :: boundary_mask_field
645 type(field_t), pointer :: node_type_field
646 type(tuple_i4_t), pointer :: marked_faces(:)
647 type(tuple_i4_t) :: marked_face
648 integer, allocatable :: dirichlet_mask_values(:)
649 integer, allocatable :: mixed_mask_values(:)
650 integer, allocatable :: resolved_mask_values(:)
651 integer :: scratch_idx(2)
652 integer :: boundary_size
653 integer :: compact_node_type_idx
654 integer :: boundary_dof_key
655 integer :: i, j, k, dof_size, m
656 integer :: dirichlet_mask_size, mixed_mask_size, resolved_mask_size
657 integer :: facet, el
658 real(kind=rp) :: bc_type
659 class(bc_t), pointer :: bc
660
661 call this%clear_masks()
662
663 call neko_scratch_registry%request_field(boundary_mask_field, &
664 scratch_idx(1), .true.)
665 call neko_scratch_registry%request_field(node_type_field, &
666 scratch_idx(2), .true.)
667
668 dof_size = this%dof%size()
669 this%face_type = 5.0_rp
670
671 ! Build a mask of all dofs on the boundary.
672 ! Scratch clearing follows the active backend, whereas the resolution
673 ! below is deliberately assembled on the host. Clear the host copy
674 ! explicitly before setting its marked entries.
675 call rzero(boundary_mask_field%x, dof_size)
676 do i = 1, this%bcs%size()
677 bc => this%bcs%get(i)
678
679 if (.not. allocated(bc%msk)) then
680 call neko_error("Attempting to finalize coupled projector " // &
681 "unfinalized BC.")
682 end if
683
684 ! Mask all the dofs touched by this BC. Since %msk is propagated to all
685 ! local dofs via gather-scatter, boundary_mask_field will contain all
686 ! local nodes on the boundary, including those elements that don't touch
687 ! it with a face.
688 ! Note that bc%msk stores its length in slot 0 and is therefore passed
689 ! with the `_0` masked-wrapper convention.
690 call cfill_mask(boundary_mask_field%x, 1.0_rp, dof_size, &
691 bc%msk(1:bc%msk(0)), &
692 bc%msk(0))
693 end do
694
695 ! Set priority values (see the BC_* constants) for constraint assignment.
696 ! Mimics the procedure in Nek5000 directly.
697 ! The values are chosen so that a min reduction applies the
698 ! highest-priority constraint.
699 ! 5 -> unconstrained
700 ! 3 -> tangentially constrained
701 ! 2 -> normally constrained
702 ! 0 -> fully constrained
703
704 ! Fill the field to not mess up gather-scatter reduction later.
705 call cfill(node_type_field%x, 5.0_rp, dof_size)
706
707 do i = 1, this%bcs%size()
708 bc => this%bcs%get(i)
709 bc_type = bc%bc_type
710
711 ! Store the type on each boundary face touched by this BC.
712 ! This is the compact analogue of Nek's face-resident HFMASK field:
713 ! one scalar type value per local (facet, element) pair.
714 marked_faces => bc%marked_facet%array()
715 do j = 1, bc%marked_facet%size()
716 marked_face = marked_faces(j)
717 facet = marked_face%x(1)
718 el = marked_face%x(2)
719 this%face_type(facet, el) = bc_type
720 end do
721
722 ! Note that facet_node_msk is used, so constraints are only directly
723 ! applied to elements that touch the boundary at this point.
724 do j = 1, bc%facet_node_msk(0)
725 m = bc%facet_node_msk(j)
726 ! The min here ensures that the highest-priority constraint is kept
727 ! within a single element.
728 node_type_field%x(m,1,1,1) = min(bc_type, node_type_field%x(m,1,1,1))
729 end do
730 end do
731
732 ! Propagate constraints to all local dofs via gather-scatter.
733 ! Ensures the highest-priority constraint is kept across element
734 ! boundaries.
735 if (neko_bcknd_device .eq. 1) then
736 call node_type_field%copy_from(host_to_device, .true.)
737 end if
738 call this%coef%gs_h%op(node_type_field, gs_op_min)
739 if (neko_bcknd_device .eq. 1) then
740 call node_type_field%copy_from(device_to_host, .true.)
741 end if
742
743 ! Build compact nodal type cache for all boundary dofs.
744 ! First pass to count the size of the boundary dof set.
745
746 boundary_size = 0
747 !$omp parallel do reduction(+:boundary_size)
748 do i = 1, dof_size
749 if (boundary_mask_field%x(i,1,1,1) .gt. 0.5_rp) then
750 boundary_size = boundary_size + 1
751 end if
752 end do
753 !$omp end parallel do
754
755 ! Linear indices of boundary dofs
756 allocate(this%boundary_dof(boundary_size))
757 ! Node BC type per boundary dof
758 allocate(this%node_type(boundary_size))
759 ! Mapping from global dof index to compact boundary dof index and type.
760 call this%boundary_idx%init(boundary_size, compact_node_type_idx)
761
762 boundary_size = 0
763 do i = 1, dof_size
764 if (boundary_mask_field%x(i,1,1,1) .lt. 0.5_rp) cycle
765
766 boundary_size = boundary_size + 1
767 this%boundary_dof(boundary_size) = i
768 this%node_type(boundary_size) = node_type_field%x(i,1,1,1)
769 boundary_dof_key = i
770 compact_node_type_idx = boundary_size
771 call this%boundary_idx%set(boundary_dof_key, compact_node_type_idx)
772 end do
773
774 ! For mixed BCs, build a resolved subset of the original bc%msk support.
775 ! A dof survives in the resolved mask only if the globally reduced type
776 ! still matches the type semantics of that BC. Nodes that were touched by
777 ! the BC originally, but whose meaning changed after shared-node reduction,
778 ! are therefore dropped here.
779 do i = 1, this%bcs%size()
780 bc => this%bcs%get(i)
781
782 select type (bc)
783 class is (mixed_bc_t)
784 call bc%resolved_msk%free()
785 call bc%n%free()
786 call bc%t1%free()
787 call bc%t2%free()
788
789 bc_type = bc%bc_type
790
791 ! First pass to count the size of the resolved mask.
792 resolved_mask_size = 0
793 do j = 1, bc%msk(0)
794 k = bc%msk(j)
795 ! Compare face and node bc type
796 if (abs(node_type_field%x(k,1,1,1) - bc_type) .lt. 1.0e-6_rp) then
797 resolved_mask_size = resolved_mask_size + 1
798 end if
799 end do
800
801 allocate(resolved_mask_values(resolved_mask_size))
802
803 ! Fill in the mask values
804 resolved_mask_size = 0
805 do j = 1, bc%msk(0)
806 k = bc%msk(j)
807 if (abs(node_type_field%x(k,1,1,1) - bc_type) .lt. 1.0e-6_rp) then
808 resolved_mask_size = resolved_mask_size + 1
809 resolved_mask_values(resolved_mask_size) = k
810 end if
811 end do
812
813 call bc%resolved_msk%init(resolved_mask_values, resolved_mask_size)
814 deallocate(resolved_mask_values)
815 end select
816 end do
817
818 ! Partition the resolved boundary dofs into the fully constrained subset
819 ! and the mixed subset. Only the latter needs a local basis.
820 dirichlet_mask_size = 0
821 mixed_mask_size = 0
822
823 !$omp parallel do reduction(+:dirichlet_mask_size,mixed_mask_size)
824 do i = 1, dof_size
825 ! Internal node
826 if (boundary_mask_field%x(i,1,1,1) .lt. 0.5_rp) cycle
827
828 if (node_type_field%x(i,1,1,1) .lt. 1.9_rp) then
829 dirichlet_mask_size = dirichlet_mask_size + 1
830 else if (node_type_field%x(i,1,1,1) .gt. 1.9_rp .and. &
831 node_type_field%x(i,1,1,1) .lt. 3.9_rp) then
832 mixed_mask_size = mixed_mask_size + 1
833 end if
834 end do
835 !$omp end parallel do
836
837 allocate(dirichlet_mask_values(dirichlet_mask_size))
838 allocate(mixed_mask_values(mixed_mask_size))
839
840 ! We reuse the variables as counters in the loop below, which actually
841 ! fills the masks, now that we know the size.
842 dirichlet_mask_size = 0
843 mixed_mask_size = 0
844 do i = 1, dof_size
845 if (boundary_mask_field%x(i,1,1,1) .lt. 0.5_rp) cycle
846
847 if (node_type_field%x(i,1,1,1) .lt. 1.9_rp) then
848 dirichlet_mask_size = dirichlet_mask_size + 1
849 dirichlet_mask_values(dirichlet_mask_size) = i
850 else if (node_type_field%x(i,1,1,1) .gt. 1.9_rp .and. &
851 node_type_field%x(i,1,1,1) .lt. 3.9_rp) then
852 mixed_mask_size = mixed_mask_size + 1
853 mixed_mask_values(mixed_mask_size) = i
854 end if
855 end do
856
857 ! Initialize both masks. The temporary arrays remain valid actual
858 ! arguments when their allocated extent is zero.
859 call this%dirichlet_dof_mask%init(dirichlet_mask_values, &
860 dirichlet_mask_size)
861 call this%mixed_dof_mask%init(mixed_mask_values, mixed_mask_size)
862
863 ! Allocate mixed-node constraints and fill them from the reduced type.
864 allocate(this%constraint_n(mixed_mask_size))
865 allocate(this%constraint_t1(mixed_mask_size))
866 allocate(this%constraint_t2(mixed_mask_size))
867 if (neko_bcknd_device .eq. 1) then
868 call device_map(this%constraint_n, this%constraint_n_d, &
869 size(this%constraint_n))
870 call device_map(this%constraint_t1, this%constraint_t1_d, &
871 size(this%constraint_t1))
872 call device_map(this%constraint_t2, this%constraint_t2_d, &
873 size(this%constraint_t2))
874 end if
875
876 !OCL NORECURRENCE, NOVREC, NOALIAS
877 !DIR$ CONCURRENT
878 !DIR$ IVDEP
879 !GCC$ ivdep
880 !$omp do
881 do i = 1, mixed_mask_size
882 j = mixed_mask_values(i)
883
884 if (node_type_field%x(j,1,1,1) .lt. 1.9_rp) then
885 this%constraint_n(i) = 1
886 this%constraint_t1(i) = 1
887 this%constraint_t2(i) = 1
888 else if (node_type_field%x(j,1,1,1) .gt. 1.9_rp .and. &
889 node_type_field%x(j,1,1,1) .lt. 2.9_rp) then
890 this%constraint_n(i) = 1
891 this%constraint_t1(i) = 0
892 this%constraint_t2(i) = 0
893 else if (node_type_field%x(j,1,1,1) .gt. 2.9_rp .and. &
894 node_type_field%x(j,1,1,1) .lt. 3.9_rp) then
895 this%constraint_n(i) = 0
896 this%constraint_t1(i) = 1
897 this%constraint_t2(i) = 1
898 else if (node_type_field%x(j,1,1,1) .gt. 3.9_rp) then
899 this%constraint_n(i) = 0
900 this%constraint_t1(i) = 0
901 this%constraint_t2(i) = 0
902 end if
903 end do
904 !$omp end do
905
906 deallocate(dirichlet_mask_values)
907 deallocate(mixed_mask_values)
908 call neko_scratch_registry%relinquish_field(scratch_idx)
910
918 class(coupled_vector_bc_projector_t), intent(inout) :: this
919 type(field_t), pointer :: normal_x_field
920 type(field_t), pointer :: normal_y_field
921 type(field_t), pointer :: normal_z_field
922 class(bc_t), pointer :: bc
923 integer, pointer :: mixed_dof_values(:)
924 integer, allocatable :: dof_to_mixed_idx(:)
925 integer :: scratch_idx(3)
926 integer :: node_type_lookup_status, compact_node_type_idx
927 integer :: i, j, k, dof_size, m
928 integer :: idx(4), facet, el, edge, node, ii, p
929 integer :: rst(3), rst1(3), rst2(3), step_rst(3)
930 integer :: edge_len, edge_idx, node_idx
931 real(kind=rp) :: normal(3), t1_vec(3), t2_vec(3), len, bc_type
932 real(kind=rp), parameter :: normal_tol = 100.0_rp * epsilon(1.0_rp)
933 character(len=LOG_SIZE) :: error_msg
934
935 call this%clear_basis()
936
937 call neko_scratch_registry%request_field(normal_x_field, scratch_idx(1), &
938 .true.)
939 call neko_scratch_registry%request_field(normal_y_field, scratch_idx(2), &
940 .true.)
941 call neko_scratch_registry%request_field(normal_z_field, scratch_idx(3), &
942 .true.)
943
944 dof_size = this%dof%size()
945 m = this%mixed_dof_mask%size()
946 mixed_dof_values => this%mixed_dof_mask%get()
947
948 call this%n%init(3, m)
949 call this%t1%init(3, m)
950 call this%t2%init(3, m)
951
952 ! A mapping between the field linear index of a mixed node into its index
953 ! in the mixed_dof_mask.
954 allocate(dof_to_mixed_idx(dof_size))
955 dof_to_mixed_idx = 0
956
957 !OCL NORECURRENCE, NOVREC, NOALIAS
958 !DIR$ CONCURRENT
959 !DIR$ IVDEP
960 !GCC$ ivdep
961 !$omp do
962 do i = 1, m
963 dof_to_mixed_idx(mixed_dof_values(i)) = i
964 end do
965 !$omp end do
966
967 ! this%face_type stores the face-based bc_type values, and this%node_type
968 ! stores them node-wise, after propagation with min reduction.
969 ! Note that the propagation means that some nodes may have a different,
970 ! higher-priority class than their owning face!
971 ! The algorithm for constructing normals below will make use of both
972 ! classifications when looking at edges and corners. We will really only
973 ! care about types 2 and 3, i.e. mixed bcs. The key question will be
974 ! whether a given face should contribute its normal to the edge and corner
975 ! dofs. The idea is that the face only contributes its normal if its type
976 ! is the same as that of the node.
977
978 ! Set normals at unambiguous boundary dofs based on face normals.
979 call rzero(normal_x_field%x, dof_size)
980 call rzero(normal_y_field%x, dof_size)
981 call rzero(normal_z_field%x, dof_size)
982
983 this%n = 0.0_rp
984 this%t1 = 0.0_rp
985 this%t2 = 0.0_rp
986
987 do i = 1, this%bcs%size()
988 bc => this%bcs%get(i)
989
990 ! First pass: seed the local normal field on all nodes that lie
991 ! on directly marked mixed faces.
992 ! Since several faces will own edge and corner nodes, the normals there
993 ! will be overwritten in arbitrary order, but we don't care because we
994 ! will reset those later and treat them specially.
995 do j = 1, bc%facet_node_msk(0)
996 ! Global linear index of the node on which to set the normal.
997 k = bc%facet_node_msk(j)
998
999 ! Grab face and ijke indices to address face_type and get_normal.
1000 facet = bc%facet(j)
1001 idx = nonlinear_index(k, this%coef%Xh%lx, this%coef%Xh%ly, &
1002 this%coef%Xh%lz)
1003
1004 if (this%face_type(facet, idx(4)) .lt. 1.9_rp .or. &
1005 this%face_type(facet, idx(4)) .gt. 3.1_rp) cycle
1006
1007 normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), &
1008 facet)
1009 normal_x_field%x(k,1,1,1) = normal(1)
1010 normal_y_field%x(k,1,1,1) = normal(2)
1011 normal_z_field%x(k,1,1,1) = normal(3)
1012 end do
1013 end do
1014
1015 !write(*,*) "Seeded normals at directly marked mixed nodes in " // &
1016 ! "coupled vector BC projector."
1017
1018 ! We now treat the special edges and conrners. Everything is done locally
1019 ! per element, using reference element address tables found in hex.f90
1020 ! and inside this type.
1021
1022 ! Mixed edge interiors are rebuilt from the normals of the adjacent
1023 ! faces whose local face type matches the reduced nodal type.
1024 ! This is the central point: if the adjacent face is a different type,
1025 ! which by construction can only be a lower-priority type, then it
1026 ! should not contribute its normal.
1027 !
1028 ! Consider the following 2D example. In 2D an edge becomes a node in the
1029 ! corner of the element. Look at the node marked with X. After the nodal
1030 ! type is propagated, it will have type 2---the highest-priority of the
1031 ! adjacent. So, only the face with type 2 in El 2 will contribute to the
1032 ! normal. This is a rather extreme example, but it illustrates well what
1033 ! can happen.
1034 !
1035 ! ---------
1036 ! | El 1 |
1037 ! | 3 |
1038 ! | |
1039 ! | 3 | 3
1040 ! --------X---------
1041 ! | El 2 |
1042 ! | |
1043 ! | 2 |
1044 ! | |
1045 ! --------0
1046
1047 ! We loop over the edges of all elements, so we catch those that touch
1048 ! the boundary with an edge or a corner but not a face. Note that we will
1049 ! only treat the interior nodes of the edge here. The endpoints, i.e.
1050 ! the corner nodes are handled in the next loop.
1051 if (m .gt. 0) then
1052 do el = 1, this%coef%msh%nelv
1053 do edge = 1, size(edge_nodes, 2)
1054 ! Representitive rst index in the middle of an edge.
1055 rst = this%edge_mid_rst(:, edge)
1056
1057 ! Global linear index of the midpoint node.
1058 edge_idx = linear_index(rst(1), rst(2), rst(3), el, &
1059 this%coef%Xh%lx, this%coef%Xh%ly, this%coef%Xh%lz)
1060
1061 ! Get the BC type.
1062 node_type_lookup_status = this%boundary_idx%get( &
1063 edge_idx, compact_node_type_idx)
1064
1065 ! If we did not find the node in the lookup, it means it is not
1066 ! a boundary node and we can skip.
1067 if (node_type_lookup_status .ne. 0) cycle
1068 bc_type = abs(this%node_type(compact_node_type_idx))
1069
1070 ! If this is not a mixed bc edge, just leave it alone.
1071 if (bc_type .lt. 1.9_rp .or. bc_type .gt. 3.1_rp) cycle
1072
1073 ! Recall that "node" in the lookup table names refer to element
1074 ! corners. This is to stay consistent with the hex_t notation.
1075
1076 ! Get edge endpoints index triples. For example,
1077 ! (1, 1, 1) and (lx, 1, 1)
1078 rst1 = this%node_rst(:, edge_nodes(1, edge))
1079 rst2 = this%node_rst(:, edge_nodes(2, edge))
1080
1081 ! Compute number of gll nodes on the edge, so lx, ly, or lz.
1082 ! Which currently in Neko is one and the same.
1083 edge_len = maxval(abs(rst2 - rst1)) + 1
1084
1085 ! The running index direction along the edge in rst-space.
1086 ! This is a triple, but only one component is nonzero.
1087 step_rst = 0
1088 do ii = 1, 3
1089 if (rst2(ii) .gt. rst1(ii)) then
1090 step_rst(ii) = 1
1091 else if (rst2(ii) .lt. rst1(ii)) then
1092 step_rst(ii) = -1
1093 end if
1094 end do
1095
1096 ! Loop over interior edge nodes and reset the normals.
1097 do p = 2, edge_len - 1
1098 rst = rst1 + (p - 1) * step_rst
1099 k = linear_index(rst(1), rst(2), rst(3), el, &
1100 this%coef%Xh%lx, this%coef%Xh%ly, this%coef%Xh%lz)
1101 normal_x_field%x(k,1,1,1) = 0.0_rp
1102 normal_y_field%x(k,1,1,1) = 0.0_rp
1103 normal_z_field%x(k,1,1,1) = 0.0_rp
1104 end do
1105
1106 ! Loop over the faces adjacent to this edge and add the normals
1107 ! if the BC type matches between the face and the edge.
1108 do ii = 1, size(edge_faces, 1)
1109 facet = edge_faces(ii, edge)
1110
1111 ! Skip if the BC type is not the same.
1112 if (abs(bc_type - this%face_type(facet, el)) .gt. 1.0e-6_rp) then
1113 cycle
1114 end if
1115
1116 ! Loop over the interior edge nodes again and add the normals.
1117 do p = 2, edge_len - 1
1118 rst = rst1 + (p - 1) * step_rst
1119 k = linear_index(rst(1), rst(2), rst(3), el, &
1120 this%coef%Xh%lx, this%coef%Xh%ly, this%coef%Xh%lz)
1121 normal = this%coef%get_normal(rst(1), rst(2), rst(3), &
1122 el, facet)
1123 normal_x_field%x(k,1,1,1) = &
1124 normal_x_field%x(k,1,1,1) + normal(1)
1125 normal_y_field%x(k,1,1,1) = &
1126 normal_y_field%x(k,1,1,1) + normal(2)
1127 normal_z_field%x(k,1,1,1) = &
1128 normal_z_field%x(k,1,1,1) + normal(3)
1129 end do
1130 end do
1131 end do
1132 end do
1133
1134 !write(*,*) "Finished reconstructing normals at mixed edges in " // &
1135 ! "coupled vector BC projector."
1136
1137 ! Mixed corner node normals are rebuilt from the adjacent faces whose
1138 ! local face type matches the reduced nodal type at that node.
1139 do el = 1, this%coef%msh%nelv
1140 do node = 1, size(this%node_linear_idx)
1141 rst = this%node_rst(:, node)
1142 node_idx = linear_index(rst(1), rst(2), rst(3), el, &
1143 this%coef%Xh%lx, this%coef%Xh%ly, this%coef%Xh%lz)
1144
1145 node_type_lookup_status = this%boundary_idx%get( &
1146 node_idx, compact_node_type_idx)
1147 if (node_type_lookup_status .ne. 0) cycle
1148 bc_type = abs(this%node_type(compact_node_type_idx))
1149
1150 ! Ignore if the BC type is not a mixed one.
1151 if (bc_type .lt. 1.9_rp .or. bc_type .gt. 3.1_rp) cycle
1152
1153 ! Kill the normal to start fresh.
1154 normal_x_field%x(node_idx,1,1,1) = 0.0_rp
1155 normal_y_field%x(node_idx,1,1,1) = 0.0_rp
1156 normal_z_field%x(node_idx,1,1,1) = 0.0_rp
1157
1158 ! Note, 3 faces share a corner node in 3D.
1159 do ii = 1, this%coef%msh%gdim
1160 facet = node_faces(ii, node)
1161
1162 ! Check type agreement
1163 if (abs(bc_type - this%face_type(facet, el)) .gt. 1.0e-6_rp) then
1164 cycle
1165 end if
1166
1167 ! Add the normal.
1168 normal = this%coef%get_normal(rst(1), rst(2), rst(3), &
1169 el, facet)
1170 normal_x_field%x(node_idx,1,1,1) = &
1171 normal_x_field%x(node_idx,1,1,1) + normal(1)
1172 normal_y_field%x(node_idx,1,1,1) = &
1173 normal_y_field%x(node_idx,1,1,1) + normal(2)
1174 normal_z_field%x(node_idx,1,1,1) = &
1175 normal_z_field%x(node_idx,1,1,1) + normal(3)
1176 end do
1177 end do
1178 end do
1179 end if
1180
1181 ! We are done element-wise. We now need global consistency at shared nodes.
1182 !
1183 ! For cyclic meshes, however, periodic counterparts may not be co-planar.
1184 ! If we directly gather Cartesian normal components, we mix vectors that are
1185 ! expressed in different local frames across the periodic map, which can
1186 ! tilt the reconstructed normal.
1187 !
1188 ! We therefore follow the same cyclic treatment used in other vector
1189 ! operators in Neko:
1190 ! 1) Rotate cyclic-marked nodes into the cyclic-normal/tangential frame.
1191 ! 2) Perform the global GS_OP_ADD while all contributions use that frame.
1192 ! 3) Rotate cyclic-marked nodes back to Cartesian.
1193 !
1194 ! This keeps MPI/shared-node averaging intact while preventing frame-mixing
1195 ! on cyclic couplings.
1196 if (neko_bcknd_device .eq. 1) then
1197 call normal_x_field%copy_from(host_to_device, .false.)
1198 call normal_y_field%copy_from(host_to_device, .false.)
1199 call normal_z_field%copy_from(host_to_device, .true.)
1200 end if
1201
1202 if (this%coef%cyclic) then
1203 call rotate_cyc(normal_x_field%x, normal_y_field%x, normal_z_field%x, &
1204 1, this%coef)
1205 end if
1206
1207 call this%coef%gs_h%op(normal_x_field, gs_op_add)
1208 call this%coef%gs_h%op(normal_y_field, gs_op_add)
1209 call this%coef%gs_h%op(normal_z_field, gs_op_add)
1210
1211 if (this%coef%cyclic) then
1212 call rotate_cyc(normal_x_field%x, normal_y_field%x, normal_z_field%x, &
1213 0, this%coef)
1214 end if
1215
1216 if (neko_bcknd_device .eq. 1) then
1217 call normal_x_field%copy_from(device_to_host, .false.)
1218 call normal_y_field%copy_from(device_to_host, .false.)
1219 call normal_z_field%copy_from(device_to_host, .true.)
1220 end if
1221
1222 ! Normalize normals and build tangential directions for all mixed nodes.
1223 ! Populate this into the basis components in the type.
1224
1225 !OCL NORECURRENCE, NOVREC, NOALIAS
1226 !DIR$ CONCURRENT
1227 !DIR$ IVDEP
1228 !GCC$ ivdep
1229 !$omp do
1230 do i = 1, m
1231 j = mixed_dof_values(i)
1232
1233 ! Normalize the normal
1234 normal(1) = normal_x_field%x(j,1,1,1)
1235 normal(2) = normal_y_field%x(j,1,1,1)
1236 normal(3) = normal_z_field%x(j,1,1,1)
1237 len = sqrt(sum(normal**2))
1238 if (len .le. normal_tol) then
1239 write(error_msg, '(A,I0,A,ES13.6,A)') &
1240 "Coupled vector BC projector could not construct a normal " // &
1241 "at local DOF ", j, " (norm = ", len, ")."
1242 call neko_error(error_msg)
1243 end if
1244
1245 this%n%x(:,i) = normal / len
1246
1247 ! Select the first tangent direction by crossing the normal with a
1248 ! coordinate axis. Use y near the z-axis, where the usual z-axis
1249 ! construction becomes ill-conditioned.
1250 if (abs(this%n%x(3,i)) .gt. 0.999_rp) then
1251 t1_vec = [ this%n%x(3,i), 0.0_rp, -this%n%x(1,i) ]
1252 else
1253 t1_vec = [ -this%n%x(2,i), this%n%x(1,i), 0.0_rp ]
1254 end if
1255 len = sqrt(sum(t1_vec**2))
1256 if (len .gt. 0.0_rp) then
1257 this%t1%x(:,i) = t1_vec / len
1258 end if
1259
1260 ! Get t2 as a cross product of n and t1.
1261 t2_vec(1) = this%n%x(2,i) * this%t1%x(3,i) - &
1262 this%n%x(3,i) * this%t1%x(2,i)
1263 t2_vec(2) = this%n%x(3,i) * this%t1%x(1,i) - &
1264 this%n%x(1,i) * this%t1%x(3,i)
1265 t2_vec(3) = this%n%x(1,i) * this%t1%x(2,i) - &
1266 this%n%x(2,i) * this%t1%x(1,i)
1267 len = sqrt(sum(t2_vec**2))
1268 if (len .gt. 0.0_rp) then
1269 this%t2%x(:,i) = t2_vec / len
1270 end if
1271 end do
1272 !$omp end do
1273
1274 if (neko_bcknd_device .eq. 1) then
1275 call device_memcpy(this%constraint_n, this%constraint_n_d, &
1276 size(this%constraint_n), host_to_device, sync = .true.)
1277 call device_memcpy(this%constraint_t1, this%constraint_t1_d, &
1278 size(this%constraint_t1), host_to_device, sync = .true.)
1279 call device_memcpy(this%constraint_t2, this%constraint_t2_d, &
1280 size(this%constraint_t2), host_to_device, sync = .true.)
1281 call this%n%copy_from(host_to_device, .true.)
1282 call this%t1%copy_from(host_to_device, .true.)
1283 call this%t2%copy_from(host_to_device, .true.)
1284 end if
1285
1286 ! Transfer the final mixed-node basis into each mixed BC on its
1287 ! resolved support, so strong application on the physical field can use
1288 ! BC-local data rather than the global projector internals.
1289 do i = 1, this%bcs%size()
1290 bc => this%bcs%get(i)
1291
1292 select type (bc)
1293 class is (mixed_bc_t)
1294 m = bc%resolved_msk%size()
1295 call bc%n%init(3, m)
1296 call bc%t1%init(3, m)
1297 call bc%t2%init(3, m)
1298
1299 do j = 1, m
1300 k = bc%resolved_msk%get(j)
1301 p = dof_to_mixed_idx(k)
1302
1303 if (p .eq. 0) then
1304 call neko_error("Mixed BC resolved_msk entry missing from " // &
1305 "the coupled projector mixed basis.")
1306 end if
1307
1308 bc%n%x(:,j) = this%n%x(:,p)
1309 bc%t1%x(:,j) = this%t1%x(:,p)
1310 bc%t2%x(:,j) = this%t2%x(:,p)
1311 end do
1312
1313 if (neko_bcknd_device .eq. 1) then
1314 call bc%n%copy_from(host_to_device, .true.)
1315 call bc%t1%copy_from(host_to_device, .true.)
1316 call bc%t2%copy_from(host_to_device, .true.)
1317 end if
1318 end select
1319 end do
1320 call neko_scratch_registry%relinquish_field(scratch_idx)
1321 if (allocated(dof_to_mixed_idx)) deallocate(dof_to_mixed_idx)
1323
1330 subroutine coupled_vector_bc_projector_apply(this, x, y, z, n, strm)
1331 class(coupled_vector_bc_projector_t), intent(in) :: this
1332 integer, intent(in) :: n
1333 real(kind=rp), intent(inout) :: x(n)
1334 real(kind=rp), intent(inout) :: y(n)
1335 real(kind=rp), intent(inout) :: z(n)
1336 type(c_ptr), intent(inout), optional :: strm
1337
1338 integer, pointer :: dirichlet_msk(:)
1339 integer, pointer :: mixed_msk(:)
1340 integer :: i, j, m
1341 real(kind=rp) :: u(3), uloc(3)
1342 type(c_ptr) :: x_d, y_d, z_d
1343 type(c_ptr):: strm_
1344
1345 if (present(strm)) then
1346 strm_ = strm
1347 else
1348 strm_ = glb_cmd_queue
1349 end if
1350
1351 ! Fully constrained nodes do not need the local basis. They are simply
1352 ! zeroed in Cartesian space before the mixed-node pass.
1353 m = this%dirichlet_dof_mask%size()
1354 if (m .gt. 0) then
1355 if (neko_bcknd_device .eq. 1) then
1356
1357
1358 x_d = device_get_ptr(x)
1359 y_d = device_get_ptr(y)
1360 z_d = device_get_ptr(z)
1361 call device_cfill_mask(x_d, 0.0_rp, n, &
1362 this%dirichlet_dof_mask%get_d(), m, strm = strm_)
1363 call device_cfill_mask(y_d, 0.0_rp, n, &
1364 this%dirichlet_dof_mask%get_d(), m, strm = strm_)
1365 call device_cfill_mask(z_d, 0.0_rp, n, &
1366 this%dirichlet_dof_mask%get_d(), m, strm = strm_)
1367 else
1368 dirichlet_msk => this%dirichlet_dof_mask%get()
1369 call cfill_mask(x, 0.0_rp, n, dirichlet_msk, m)
1370 call cfill_mask(y, 0.0_rp, n, dirichlet_msk, m)
1371 call cfill_mask(z, 0.0_rp, n, dirichlet_msk, m)
1372 end if
1373 end if
1374
1375 m = this%mixed_dof_mask%size()
1376
1377 if (m .gt. 0) then
1378 if (neko_bcknd_device .eq. 1) then
1379 x_d = device_get_ptr(x)
1380 y_d = device_get_ptr(y)
1381 z_d = device_get_ptr(z)
1383 this%mixed_dof_mask%get_d(), x_d, y_d, z_d, &
1384 this%constraint_n_d, this%constraint_t1_d, &
1385 this%constraint_t2_d, this%n%x_d, this%t1%x_d, &
1386 this%t2%x_d, m, strm_)
1387 else
1388 mixed_msk => this%mixed_dof_mask%get()
1389
1390 do i = 1, m
1391 j = mixed_msk(i)
1392
1393 u(1) = x(j)
1394 u(2) = y(j)
1395 u(3) = z(j)
1396
1397 uloc(1) = u(1) * this%n%x(1,i) + u(2) * this%n%x(2,i) + &
1398 u(3) * this%n%x(3,i)
1399 uloc(2) = u(1) * this%t1%x(1,i) + u(2) * this%t1%x(2,i) + &
1400 u(3) * this%t1%x(3,i)
1401 uloc(3) = u(1) * this%t2%x(1,i) + u(2) * this%t2%x(2,i) + &
1402 u(3) * this%t2%x(3,i)
1403
1404 if (this%constraint_n(i) .ne. 0) uloc(1) = 0.0_rp
1405 if (this%constraint_t1(i) .ne. 0) uloc(2) = 0.0_rp
1406 if (this%constraint_t2(i) .ne. 0) uloc(3) = 0.0_rp
1407
1408 u = uloc(1) * this%n%x(:,i) + uloc(2) * this%t1%x(:,i) + &
1409 uloc(3) * this%t2%x(:,i)
1410
1411 x(j) = u(1)
1412 y(j) = u(2)
1413 z(j) = u(3)
1414 end do
1415 end if
1416 end if
1418
1422 subroutine coupled_vector_bc_projector_debug_output(this, field_name)
1424 class(coupled_vector_bc_projector_t), intent(inout) :: this
1425 character(len=*), intent(in), optional :: field_name
1426 type(field_t), pointer :: mask_field
1427 type(field_t), pointer :: nx_field, ny_field, nz_field
1428 type(field_list_t) :: basis_fields
1429 type(fld_file_t) :: basis_file
1430 integer :: scratch_idx(4)
1431 integer, pointer :: mixed_mask_values(:)
1432 integer :: dof_size, mixed_mask_size
1433 character(len=:), allocatable :: field_name_
1434
1435 if (present(field_name)) then
1436 field_name_ = trim(field_name)
1437 else
1438 field_name_ = 'bc_projector'
1439 end if
1440
1441 call neko_scratch_registry%request_field(mask_field, scratch_idx(1), .true.)
1442 call neko_scratch_registry%request_field(nx_field, scratch_idx(2), .true.)
1443 call neko_scratch_registry%request_field(ny_field, scratch_idx(3), .true.)
1444 call neko_scratch_registry%request_field(nz_field, scratch_idx(4), .true.)
1445
1446 dof_size = this%dof%size()
1447 mixed_mask_size = this%mixed_dof_mask%size()
1448
1449 call rzero(mask_field%x, dof_size)
1450 call rzero(nx_field%x, dof_size)
1451 call rzero(ny_field%x, dof_size)
1452 call rzero(nz_field%x, dof_size)
1453
1454 if (this%mixed_dof_mask%is_set()) then
1455 mixed_mask_values => this%mixed_dof_mask%get()
1456 call masked_scatter_copy(nx_field%x(:,1,1,1), this%n%x(1,:), &
1457 mixed_mask_values, dof_size, mixed_mask_size)
1458 call masked_scatter_copy(ny_field%x(:,1,1,1), this%n%x(2,:), &
1459 mixed_mask_values, dof_size, mixed_mask_size)
1460 call masked_scatter_copy(nz_field%x(:,1,1,1), this%n%x(3,:), &
1461 mixed_mask_values, dof_size, mixed_mask_size)
1462 end if
1463
1464 if (neko_bcknd_device .eq. 1) then
1465 call device_cfill(mask_field%x_d, 5.0_rp, dof_size)
1466
1467 if (this%dirichlet_dof_mask%is_set()) then
1468 call device_cfill_mask(mask_field%x_d, 1.0_rp, dof_size, &
1469 this%dirichlet_dof_mask%get_d(), &
1470 this%dirichlet_dof_mask%size())
1471 end if
1472
1473 if (this%mixed_dof_mask%is_set()) then
1474 call device_cfill_mask(mask_field%x_d, 2.0_rp, dof_size, &
1475 this%mixed_dof_mask%get_d(), this%mixed_dof_mask%size())
1476 end if
1477
1478 call device_memcpy(mask_field%x, mask_field%x_d, dof_size, &
1480 sync = .true.)
1481 else
1482 call cfill(mask_field%x, 5.0_rp, dof_size)
1483
1484 if (this%dirichlet_dof_mask%is_set()) then
1485 call cfill_mask(mask_field%x, 1.0_rp, dof_size, &
1486 this%dirichlet_dof_mask%get(), this%dirichlet_dof_mask%size())
1487 end if
1488
1489 if (this%mixed_dof_mask%is_set()) then
1490 call cfill_mask(mask_field%x, 2.0_rp, dof_size, &
1491 this%mixed_dof_mask%get(), this%mixed_dof_mask%size())
1492 end if
1493 end if
1494
1495 call basis_fields%init(4)
1496 call basis_fields%assign(1, mask_field)
1497 call basis_fields%assign(2, nx_field)
1498 call basis_fields%assign(3, ny_field)
1499 call basis_fields%assign(4, nz_field)
1500
1501 call basis_file%init(field_name_ // '.fld')
1502 call basis_file%write(basis_fields)
1503 call basis_fields%free()
1504
1505 call neko_scratch_registry%relinquish_field(scratch_idx)
1507
1519 this, x, y, z, n, field_name)
1520 class(coupled_vector_bc_projector_t), intent(inout) :: this
1521 integer, intent(in) :: n
1522 real(kind=rp), intent(in) :: x(n)
1523 real(kind=rp), intent(in) :: y(n)
1524 real(kind=rp), intent(in) :: z(n)
1525 character(len=*), intent(in), optional :: field_name
1526 type(field_t), pointer :: normal_component_field
1527 type(field_t), pointer :: normal_component_coef_field
1528 type(field_list_t) :: output_fields
1529 type(fld_file_t) :: output_file
1530 type(field_t), pointer :: projector_nx_field, projector_ny_field, projector_nz_field
1531 type(field_t), pointer :: coef_nx_field, coef_ny_field, coef_nz_field
1532 type(field_list_t) :: normals_fields
1533 type(fld_file_t) :: normals_file
1534 integer :: scratch_idx(8)
1535 integer, pointer :: mixed_mask_values(:)
1536 integer, allocatable :: mixed_lut(:)
1537 integer :: dof_size, mixed_mask_size
1538 integer :: i, j, m, k, facet
1539 integer :: idx(4)
1540 real(kind=rp) :: coef_normal(3)
1541 class(bc_t), pointer :: bc
1542 character(len=:), allocatable :: field_name_
1543
1544 if (present(field_name)) then
1545 field_name_ = trim(field_name)
1546 else
1547 field_name_ = 'bc_projector_normal_component'
1548 end if
1549
1550 call neko_scratch_registry%request_field(normal_component_field, &
1551 scratch_idx(1), .true.)
1552 call neko_scratch_registry%request_field(normal_component_coef_field, &
1553 scratch_idx(2), .true.)
1554 call neko_scratch_registry%request_field(projector_nx_field, &
1555 scratch_idx(3), .true.)
1556 call neko_scratch_registry%request_field(projector_ny_field, &
1557 scratch_idx(4), .true.)
1558 call neko_scratch_registry%request_field(projector_nz_field, &
1559 scratch_idx(5), .true.)
1560 call neko_scratch_registry%request_field(coef_nx_field, &
1561 scratch_idx(6), .true.)
1562 call neko_scratch_registry%request_field(coef_ny_field, &
1563 scratch_idx(7), .true.)
1564 call neko_scratch_registry%request_field(coef_nz_field, &
1565 scratch_idx(8), .true.)
1566
1567 dof_size = this%dof%size()
1568 mixed_mask_size = this%mixed_dof_mask%size()
1569 call rzero(normal_component_field%x, dof_size)
1570 call rzero(normal_component_coef_field%x, dof_size)
1571 call rzero(projector_nx_field%x, dof_size)
1572 call rzero(projector_ny_field%x, dof_size)
1573 call rzero(projector_nz_field%x, dof_size)
1574 call rzero(coef_nx_field%x, dof_size)
1575 call rzero(coef_ny_field%x, dof_size)
1576 call rzero(coef_nz_field%x, dof_size)
1577
1578 allocate(mixed_lut(dof_size))
1579 mixed_lut = 0
1580 if (this%mixed_dof_mask%is_set()) then
1581 mixed_mask_values => this%mixed_dof_mask%get()
1582 do i = 1, mixed_mask_size
1583 mixed_lut(mixed_mask_values(i)) = i
1584 end do
1585 end if
1586
1587 do i = 1, this%bcs%size()
1588 bc => this%bcs%get(i)
1589 do m = 1, bc%facet_node_msk(0)
1590 k = bc%facet_node_msk(m)
1591 facet = bc%facet(m)
1592 idx = nonlinear_index(k, this%coef%Xh%lx, this%coef%Xh%ly, &
1593 this%coef%Xh%lz)
1594
1595 coef_normal = this%coef%get_normal(idx(1), idx(2), idx(3), idx(4), &
1596 facet)
1597 coef_nx_field%x(k,1,1,1) = coef_normal(1)
1598 coef_ny_field%x(k,1,1,1) = coef_normal(2)
1599 coef_nz_field%x(k,1,1,1) = coef_normal(3)
1600 normal_component_coef_field%x(k,1,1,1) = &
1601 x(k) * coef_normal(1) + y(k) * coef_normal(2) + z(k) * coef_normal(3)
1602
1603 j = mixed_lut(k)
1604 if (j .gt. 0) then
1605 projector_nx_field%x(k,1,1,1) = this%n%x(1,j)
1606 projector_ny_field%x(k,1,1,1) = this%n%x(2,j)
1607 projector_nz_field%x(k,1,1,1) = this%n%x(3,j)
1608 normal_component_field%x(k,1,1,1) = &
1609 x(k) * this%n%x(1,j) + y(k) * this%n%x(2,j) + z(k) * this%n%x(3,j)
1610 end if
1611 end do
1612 end do
1613
1614 deallocate(mixed_lut)
1615
1616 call output_fields%init(2)
1617 call output_fields%assign(1, normal_component_field)
1618 call output_fields%assign(2, normal_component_coef_field)
1619 call output_file%init(field_name_ // '.fld')
1620 call output_file%write(output_fields)
1621 call output_fields%free()
1622
1623 call normals_fields%init(3)
1624 call normals_fields%assign(1, projector_nx_field)
1625 call normals_fields%assign(2, projector_ny_field)
1626 call normals_fields%assign(3, projector_nz_field)
1627 call normals_file%init(field_name_ // '_projector_normals.fld')
1628 call normals_file%write(normals_fields)
1629 call normals_fields%free()
1630
1631 call normals_fields%init(3)
1632 call normals_fields%assign(1, coef_nx_field)
1633 call normals_fields%assign(2, coef_ny_field)
1634 call normals_fields%assign(3, coef_nz_field)
1635 call normals_file%init(field_name_ // '_coef_normals.fld')
1636 call normals_file%write(normals_fields)
1637 call normals_fields%free()
1638
1639 call neko_scratch_registry%relinquish_field(scratch_idx)
1641
1642end module vector_bc_projector
__inline__ __device__ void nonlinear_index(const int idx, const int lx, int *index)
Definition bc_utils.h:44
Return the device pointer for an associated Fortran array.
Definition device.F90:113
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
Apply cyclic boundary condition to a vector field.
Apply the resolved vector boundary constraints.
Finalize the vector boundary-condition projector.
Free the vector boundary-condition projector.
Initialize the vector boundary-condition projector.
Mark a boundary condition in the vector boundary-condition projector.
Mark a list of boundary conditions in the projector.
Defines a list of bc_t.
Definition bc_list.f90:34
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_coupled_vector_bc_projector_apply(mixed_msk, x, y, z, constraint_n, constraint_t1, constraint_t2, n, t1, t2, m, strm)
subroutine, public device_cfill_mask(a_d, c, n, mask_d, n_mask, strm)
Fill a constant to a masked vector. .
subroutine, public device_cfill(a_d, c, n, strm)
Set all elements to a constant c .
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
Defines a mapping of the degrees of freedom.
Definition dofmap.f90:35
Defines a field.
Definition field.f90:34
NEKTON fld file format.
Definition fld_file.f90:35
Defines Gather-scatter operations.
Definition gs_ops.f90:34
integer, parameter, public gs_op_add
Definition gs_ops.f90:36
integer, parameter, public gs_op_min
Definition gs_ops.f90:36
Defines a hexahedron element.
Definition hex.f90:34
integer, dimension(2, 12), parameter, public edge_faces
Face ids incident to each edge.
Definition hex.f90:158
integer, dimension(3, 8), parameter, public node_faces
Face ids incident to each node.
Definition hex.f90:136
integer, dimension(2, 12), parameter, public edge_nodes
Edge node ids.
Definition hex.f90:120
Implements a hash table ADT.
Definition htable.f90:52
Logging routines.
Definition log.f90:34
integer, parameter, public log_size
Definition log.f90:46
Object for handling masks in Neko.
Definition mask.f90:34
Definition math.f90:60
subroutine, public masked_scatter_copy(a, b, mask, n, n_mask)
Scatter a contigous vector to masked positions in a target array .
Definition math.f90:473
subroutine, public cfill(a, c, n)
Set all elements to a constant c .
Definition math.f90:600
subroutine, public rzero(a, n)
Zero a real vector.
Definition math.f90:238
subroutine, public cfill_mask(a, c, n, mask, n_mask)
Fill a constant to a masked vector. .
Definition math.f90:491
Defines a matrix.
Definition matrix.f90:34
Implements mixed_bc_t.
Definition mixed_bc.f90:31
Build configurations.
integer, parameter neko_bcknd_device
integer, parameter, public rp
Global precision used in computations.
Definition num_types.f90:14
Operators.
Definition operators.f90:34
Implements scalar_projector_t.
Defines a registry for storing and requesting temporary objects This can be used when you have a func...
type(scratch_registry_t), target, public neko_scratch_registry
Global scratch registry.
Implements a n-tuple.
Definition tuple.f90:41
Utilities.
Definition utils.f90:35
pure integer function, public linear_index(i, j, k, l, lx, ly, lz)
Compute the address of a (i,j,k,l) array with sizes (1:lx, 1:ly, 1:lz, :)
Definition utils.f90:289
Implements boundary condition projectors for vector fields. Two types concrete types are provided: se...
subroutine coupled_vector_bc_projector_mark_bc_list(this, bclst)
Register a list of vector boundary conditions in the coupled projector.
subroutine segregated_vector_bc_projector_finalize(this, rebuild_mask)
Finalize the segregated vector boundary-condition projector.
subroutine vector_bc_projector_mark_bc_list_component(this, bclst, component)
Report unsupported component-specific list marking.
subroutine segregated_vector_bc_projector_init(this, coef)
Constructor.
subroutine, public vector_bc_projector_components(this, x, y, z)
Access the component scalar projectors from a segregated vector projector.
subroutine coupled_vector_bc_projector_clear_masks(this)
Clear the resolved mask-side state of the coupled projector.
subroutine coupled_vector_bc_projector_debug_output(this, field_name)
Write fields showing the coupled projector mask and basis.
subroutine coupled_vector_bc_projector_apply(this, x, y, z, n, strm)
Apply homogeneous boundary constraints in the local basis.
subroutine coupled_vector_bc_projector_debug_output_normal_component(this, x, y, z, n, field_name)
Write scalar fields with the normal component u.n on mixed BC nodes.
subroutine segregated_vector_bc_projector_free(this)
Destructor.
subroutine segregated_vector_bc_projector_mark_bc(this, bc)
Mark a boundary condition in the segregated projector.
subroutine coupled_vector_bc_projector_clear_basis(this)
Clear the resolved basis-side state of the coupled projector.
subroutine segregated_vector_bc_projector_mark_bc_list_component(this, bclst, component)
Mark a list of boundary conditions on one Cartesian component.
subroutine coupled_vector_bc_projector_init(this, coef)
Constructor.
subroutine vector_bc_projector_mark_bc_component(this, bc, component)
Report unsupported component-specific marking.
subroutine segregated_vector_bc_projector_apply(this, x, y, z, n, strm)
Apply the segregated vector boundary constraints.
subroutine segregated_vector_bc_projector_mark_bc_component(this, bc, component)
Mark a boundary condition on one Cartesian component.
subroutine coupled_vector_bc_projector_rebuild_masks(this)
Rebuild the resolved masks and local constraint flags.
subroutine coupled_vector_bc_projector_free(this)
Destructor.
subroutine coupled_vector_bc_projector_rebuild_basis(this)
Rebuild the local mixed-node basis.
subroutine coupled_vector_bc_projector_finalize(this, rebuild_mask)
Finalize the coupled projector by resolving the accumulated BC list.
subroutine coupled_vector_bc_projector_mark_bc(this, bc)
Register one vector boundary condition in the coupled projector.
subroutine segregated_vector_bc_projector_mark_bc_list(this, bclst)
Mark a list of boundary conditions in the segregated projector.
Base type for a boundary condition.
Definition bc.f90:72
A list of allocatable `bc_t`. Follows the standard interface of lists.
Definition bc_list.f90:49
Coefficients defined on a given (mesh, ) tuple. Arrays use indices (i,j,k,e): element e,...
Definition coef.f90:93
field_list_t, To be able to group fields together
Interface for NEKTON fld files.
Definition fld_file.f90:66
Integer based hash table.
Definition htable.f90:102
Type for consistently handling masks in Neko. This type encapsulates the mask array and its associate...
Definition mask.f90:51
Base type for mixed boundary conditions that need projector-provided local-basis data on the physical...
Definition mixed_bc.f90:50
Projector for scalar boundary conditions.
Integer based 2-tuple.
Definition tuple.f90:58
A coupled projector for vector fields, suitable for mixed boundary conditions.
A projector for vector fields that acts component-wise.
Abstract type for resolving vector boundary conditions.