! //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// ! ! ! Maintainers : support@fluidnumerics.com ! Official Repository : https://github.com/FluidNumerics/self/ ! ! Copyright © 2024 Fluid Numerics LLC ! ! Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: ! ! 1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. ! ! 2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in ! the documentation and/or other materials provided with the distribution. ! ! 3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from ! this software without specific prior written permission. ! ! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS “AS IS” AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT ! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT ! HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT ! LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY ! THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF ! THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. ! ! //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// ! module SELF_RefinementIndicator_2D use SELF_Constants use SELF_RefinementIndicator_2D_t use SELF_Scalar_2D use SELF_GPU use SELF_GPUInterfaces use iso_c_binding implicit none type,extends(RefinementIndicator2D_t),public :: RefinementIndicator2D character(3) :: backend = "gpu" type(c_ptr) :: Pmodal_gpu = c_null_ptr type(c_ptr) :: indicator_gpu = c_null_ptr type(c_ptr) :: flag_gpu = c_null_ptr type(c_ptr) :: gate_gpu = c_null_ptr type(c_ptr) :: energyWeight_gpu = c_null_ptr integer :: nVarWeights_gpu = 0 !! Allocated length of energyWeight_gpu. Zero until the first Estimate fixes the solution's !! variable count, which is not known at Init. contains procedure,public :: Init => Init_RefinementIndicator2D procedure,public :: Free => Free_RefinementIndicator2D procedure,public :: UpdateHost => UpdateHost_RefinementIndicator2D procedure,public :: UpdateDevice => UpdateDevice_RefinementIndicator2D procedure,public :: Estimate => Estimate_RefinementIndicator2D endtype RefinementIndicator2D contains subroutine Init_RefinementIndicator2D(this,interp,nElem,refineThreshold,coarsenThreshold) implicit none class(RefinementIndicator2D),intent(out) :: this type(Lagrange),intent(in),target :: interp integer,intent(in) :: nElem real(prec),intent(in) :: refineThreshold real(prec),intent(in) :: coarsenThreshold ! Reuse the portable host initialization (allocation + modal transform build). call Init_RefinementIndicator2D_t(this,interp,nElem,refineThreshold,coarsenThreshold) ! The device kernel uses per-thread scratch sized to (N+1) <= 16 (AMR2D_MAXNP in ! SELF_Refinement.cpp). Fail early rather than overrun that bound on the device. if(interp%N+1 > 16) then print*,__FILE__,':',__LINE__, & ' : Error : GPU refinement indicator supports degree N <= 15 '// & '(raise AMR2D_MAXNP in SELF_Refinement.cpp to extend).' stop 1 endif call gpuCheck(hipMalloc(this%Pmodal_gpu,sizeof(this%Pmodal))) call gpuCheck(hipMalloc(this%indicator_gpu,sizeof(this%indicator))) call gpuCheck(hipMalloc(this%flag_gpu,sizeof(this%flag))) call gpuCheck(hipMalloc(this%gate_gpu,sizeof(this%gate))) ! The gate weights are sized to the solution's variable count, which the first Estimate ! establishes; nothing is allocated for them here. this%energyWeight_gpu = c_null_ptr this%nVarWeights_gpu = 0 call this%UpdateDevice() endsubroutine Init_RefinementIndicator2D subroutine Free_RefinementIndicator2D(this) implicit none class(RefinementIndicator2D),intent(inout) :: this if(c_associated(this%Pmodal_gpu)) call gpuCheck(hipFree(this%Pmodal_gpu)) if(c_associated(this%indicator_gpu)) call gpuCheck(hipFree(this%indicator_gpu)) if(c_associated(this%flag_gpu)) call gpuCheck(hipFree(this%flag_gpu)) if(c_associated(this%gate_gpu)) call gpuCheck(hipFree(this%gate_gpu)) if(c_associated(this%energyWeight_gpu)) call gpuCheck(hipFree(this%energyWeight_gpu)) this%Pmodal_gpu = c_null_ptr this%indicator_gpu = c_null_ptr this%flag_gpu = c_null_ptr this%gate_gpu = c_null_ptr this%energyWeight_gpu = c_null_ptr this%nVarWeights_gpu = 0 call Free_RefinementIndicator2D_t(this) endsubroutine Free_RefinementIndicator2D subroutine UpdateHost_RefinementIndicator2D(this) implicit none class(RefinementIndicator2D),intent(inout) :: this call gpuCheck(hipMemcpy(c_loc(this%indicator),this%indicator_gpu, & sizeof(this%indicator),hipMemcpyDeviceToHost)) call gpuCheck(hipMemcpy(c_loc(this%flag),this%flag_gpu, & sizeof(this%flag),hipMemcpyDeviceToHost)) call gpuCheck(hipMemcpy(c_loc(this%gate),this%gate_gpu, & sizeof(this%gate),hipMemcpyDeviceToHost)) endsubroutine UpdateHost_RefinementIndicator2D subroutine UpdateDevice_RefinementIndicator2D(this) implicit none class(RefinementIndicator2D),intent(inout) :: this call gpuCheck(hipMemcpy(this%Pmodal_gpu,c_loc(this%Pmodal), & sizeof(this%Pmodal),hipMemcpyHostToDevice)) call gpuCheck(hipMemcpy(this%indicator_gpu,c_loc(this%indicator), & sizeof(this%indicator),hipMemcpyHostToDevice)) call gpuCheck(hipMemcpy(this%flag_gpu,c_loc(this%flag), & sizeof(this%flag),hipMemcpyHostToDevice)) call gpuCheck(hipMemcpy(this%gate_gpu,c_loc(this%gate), & sizeof(this%gate),hipMemcpyHostToDevice)) endsubroutine UpdateDevice_RefinementIndicator2D 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 endmodule SELF_RefinementIndicator_2D