Estimate_RefinementIndicator2D Subroutine

  1. SELF_RefinementIndicator_2D.f90
  2. SELF_RefinementIndicator_2D
  3. Estimate_RefinementIndicator2D

public subroutine Estimate_RefinementIndicator2D(this, solution, ivar, comm, gate)

GPU path: one device thread per element computes the modal transform, the raw smoothness ratio and the gate energy; those are copied back to the host, where the shared second phase applies the amplitude gate, the log10 and the thresholds.

The second phase is host-side (rather than a second device kernel) because the amplitude gate needs a field-wide energy scale - a reduction, and an MPI collective when the mesh is decomposed - that an element-local device thread cannot supply. Running it through the same FinalizeIndicator the portable backend uses also makes the flags identical on CPU and GPU by construction. The extra traffic is one per-element array each way, once per adaptation epoch; nothing is added to the time-stepping loop.

Arguments

TypeIntentOptionalAttributesName
class(RefinementIndicator2D), intent(inout) :: this
class(Scalar2D), intent(in) :: solution
integer, intent(in) :: ivar
integer, intent(in), optional :: comm
real(kind=prec), intent(in), optional :: gate(:)

Calls

proc~~estimate_refinementindicator2d~~CallsGraph proc~estimate_refinementindicator2d Estimate_RefinementIndicator2D proc~checkestimatearguments~2 CheckEstimateArguments proc~estimate_refinementindicator2d->proc~checkestimatearguments~2 proc~resolveenergyweights~2 ResolveEnergyWeights proc~estimate_refinementindicator2d->proc~resolveenergyweights~2 interface~hipmalloc hipMalloc proc~estimate_refinementindicator2d->interface~hipmalloc proc~gpucheck gpuCheck proc~estimate_refinementindicator2d->proc~gpucheck interface~hipfree hipFree proc~estimate_refinementindicator2d->interface~hipfree interface~hipmemcpy hipMemcpy proc~estimate_refinementindicator2d->interface~hipmemcpy interface~refinementindicator_2d_gpu RefinementIndicator_2D_gpu proc~estimate_refinementindicator2d->interface~refinementindicator_2d_gpu proc~resolveenergyscale~2 ResolveEnergyScale proc~estimate_refinementindicator2d->proc~resolveenergyscale~2 proc~finalizeindicator~2 FinalizeIndicator proc~estimate_refinementindicator2d->proc~finalizeindicator~2 mpi_allreduce mpi_allreduce proc~resolveenergyscale~2->mpi_allreduce

Contents

Source Code

Estimate_RefinementIndicator2D

Source Code

  subroutine Estimate_RefinementIndicator2D(this,solution,ivar,comm,gate)
    !! GPU path: one device thread per element computes the modal transform, the raw smoothness
    !! ratio and the gate energy; those are copied back to the host, where the shared second phase
    !! applies the amplitude gate, the log10 and the thresholds.
    !!
    !! The second phase is host-side (rather than a second device kernel) because the amplitude
    !! gate needs a field-wide energy scale - a reduction, and an MPI collective when the mesh is
    !! decomposed - that an element-local device thread cannot supply. Running it through the same
    !! FinalizeIndicator the portable backend uses also makes the flags identical on CPU and GPU
    !! by construction. The extra traffic is one per-element array each way, once per adaptation
    !! epoch; nothing is added to the time-stepping loop.
    implicit none
    class(RefinementIndicator2D),intent(inout) :: this
    class(Scalar2D),intent(in) :: solution
    integer,intent(in) :: ivar
    integer,intent(in),optional :: comm
    real(prec),intent(in),optional :: gate(:)
    ! Local
    integer :: iel
    real(prec) :: energyScale

    call CheckEstimateArguments(this,solution,ivar,gate)
    call ResolveEnergyWeights(this,solution%nVar,ivar)

    if(this%nVarWeights_gpu /= this%nVarWeights) then
      if(c_associated(this%energyWeight_gpu)) call gpuCheck(hipFree(this%energyWeight_gpu))
      call gpuCheck(hipMalloc(this%energyWeight_gpu,sizeof(this%energyWeight)))
      this%nVarWeights_gpu = this%nVarWeights
    endif
    call gpuCheck(hipMemcpy(this%energyWeight_gpu,c_loc(this%energyWeight), &
                            sizeof(this%energyWeight),hipMemcpyHostToDevice))

    ! First phase on the device. indicator_gpu transiently holds the RAW ratio S_e, not log10(S_e).
    call RefinementIndicator_2D_gpu(this%Pmodal_gpu,solution%interior_gpu, &
                                    this%energyWeight_gpu, &
                                    this%indicator_gpu,this%gate_gpu, &
                                    this%N,solution%nVar,ivar,this%nElem)

    call this%UpdateHost()

    if(present(gate)) then
      do iel = 1,this%nElem
        this%gate(iel) = gate(iel)
      enddo
    endif

    call ResolveEnergyScale(this,energyScale,comm)
    call FinalizeIndicator(this,energyScale)

    ! Re-establish the device mirrors of indicator/flag/gate, which the host phase has just
    ! rewritten, so the device copies are never left stale.
    call this%UpdateDevice()

  endsubroutine Estimate_RefinementIndicator2D