GPU regrid: release the device copies of the old boundary-condition element/side arrays, rebuild the model storage and BC maps on the new mesh (base Regrid), then upload the new BC arrays to the device - the same device bookkeeping Init/Free do around the base Init/Free.
| Type | Intent | Optional | Attributes | Name | ||
|---|---|---|---|---|---|---|
| class(DGModel2D), | intent(inout) | :: | this | |||
| type(Mesh2D), | intent(in), | target | :: | mesh | ||
| type(SEMQuad), | intent(in), | target | :: | geometry |
subroutine Regrid_DGModel2D(this,mesh,geometry)
!! GPU regrid: release the device copies of the old boundary-condition element/side
!! arrays, rebuild the model storage and BC maps on the new mesh (base Regrid), then
!! upload the new BC arrays to the device - the same device bookkeeping Init/Free do
!! around the base Init/Free.
implicit none
class(DGModel2D),intent(inout) :: this
type(Mesh2D),intent(in),target :: mesh
type(SEMQuad),intent(in),target :: geometry
! Local
type(BoundaryCondition),pointer :: bc
! Free hyperbolic BC device arrays (old mesh)
bc => this%hyperbolicBCs%head
do while(associated(bc))
if(c_associated(bc%elements_gpu)) call gpuCheck(hipFree(bc%elements_gpu))
if(c_associated(bc%sides_gpu)) call gpuCheck(hipFree(bc%sides_gpu))
bc%elements_gpu = c_null_ptr
bc%sides_gpu = c_null_ptr
bc => bc%next
enddo
! Free parabolic BC device arrays (old mesh)
bc => this%parabolicBCs%head
do while(associated(bc))
if(c_associated(bc%elements_gpu)) call gpuCheck(hipFree(bc%elements_gpu))
if(c_associated(bc%sides_gpu)) call gpuCheck(hipFree(bc%sides_gpu))
bc%elements_gpu = c_null_ptr
bc%sides_gpu = c_null_ptr
bc => bc%next
enddo
call Regrid_DGModel2D_t(this,mesh,geometry)
! Upload hyperbolic BC element/side arrays to device (new mesh)
bc => this%hyperbolicBCs%head
do while(associated(bc))
if(bc%nBoundaries > 0) then
call gpuCheck(hipMalloc(bc%elements_gpu,sizeof(bc%elements)))
call gpuCheck(hipMemcpy(bc%elements_gpu,c_loc(bc%elements), &
sizeof(bc%elements),hipMemcpyHostToDevice))
call gpuCheck(hipMalloc(bc%sides_gpu,sizeof(bc%sides)))
call gpuCheck(hipMemcpy(bc%sides_gpu,c_loc(bc%sides), &
sizeof(bc%sides),hipMemcpyHostToDevice))
endif
bc => bc%next
enddo
! Upload parabolic BC element/side arrays to device (new mesh)
bc => this%parabolicBCs%head
do while(associated(bc))
if(bc%nBoundaries > 0) then
call gpuCheck(hipMalloc(bc%elements_gpu,sizeof(bc%elements)))
call gpuCheck(hipMemcpy(bc%elements_gpu,c_loc(bc%elements), &
sizeof(bc%elements),hipMemcpyHostToDevice))
call gpuCheck(hipMalloc(bc%sides_gpu,sizeof(bc%sides)))
call gpuCheck(hipMemcpy(bc%sides_gpu,c_loc(bc%sides), &
sizeof(bc%sides),hipMemcpyHostToDevice))
endif
bc => bc%next
enddo
endsubroutine Regrid_DGModel2D