delete scaling CF/FC advection diffusion

src/gpu/core/Calculation/UpdateGrid27.cpp

@@ -316,28 +316,11 @@ void UpdateGrid27::preCollisionBC(int level, unsigned int t)
 void UpdateGrid27::fineToCoarse(int level, InterpolationCells* fineToCoarse, ICellNeigh &neighborFineToCoarse, CudaStreamIndex streamIndex)
 {
     gridScalingKernelManager->runFineToCoarseKernelLB(level, fineToCoarse, neighborFineToCoarse, streamIndex);
-
-    if (para->getDiffOn()) {
-        if (para->getStreamManager()->streamIsRegistered(streamIndex)) {
-            printf("fineToCoarse Advection Diffusion not implemented"); // TODO
-            return;
-        }
-        gridScalingKernelManager->runFineToCoarseKernelAD(level);
-    }
 }
 
 void UpdateGrid27::coarseToFine(int level, InterpolationCells* coarseToFine, ICellNeigh &neighborCoarseToFine, CudaStreamIndex streamIndex)
 {
     this->gridScalingKernelManager->runCoarseToFineKernelLB(level, coarseToFine, neighborCoarseToFine, streamIndex);
-
-    if (para->getDiffOn())
-    {
-        if(para->getStreamManager()->streamIsRegistered(streamIndex)){
-            printf("CoarseToFineWithStream Advection Diffusion not implemented"); // TODO
-            return;
-        }
-        this->gridScalingKernelManager->runCoarseToFineKernelAD(level);
-    }
 }
 
 void UpdateGrid27::interactWithActuators(int level, unsigned int t)

src/gpu/core/GridScaling/GridScalingKernelManager.cpp

@@ -64,58 +64,8 @@ void GridScalingKernelManager::runFineToCoarseKernelLB(const int level, Interpol
     this->scalingFineToCoarse(para->getParD(level).get(), para->getParD(level+1).get(), fineToCoarse, neighborFineToCoarse, stream);
 }
 
-void GridScalingKernelManager::runFineToCoarseKernelAD(const int level) const
-{
-    scaleFineToCoarseAdvectionDiffusion(
-        para->getParD(level)->distributions.f[0],
-        para->getParD(level+1)->distributions.f[0],
-        para->getParD(level)->distributionsAD.f[0],
-        para->getParD(level+1)->distributionsAD.f[0],
-        para->getParD(level)->neighborX,
-        para->getParD(level)->neighborY,
-        para->getParD(level)->neighborZ,
-        para->getParD(level+1)->neighborX,
-        para->getParD(level+1)->neighborY,
-        para->getParD(level+1)->neighborZ,
-        para->getParD(level)->numberOfNodes,
-        para->getParD(level+1)->numberOfNodes,
-        para->getParD(level)->isEvenTimestep,
-        para->getParD(level)->fineToCoarse.coarseCellIndices,
-        para->getParD(level)->fineToCoarse.fineCellIndices,
-        para->getParD(level)->fineToCoarse.numberOfCells,
-        para->getParD(level)->viscosity,
-        para->getParD(level)->diffusivity,
-        para->getParD(level)->numberofthreads,
-        para->getParD(level)->neighborFineToCoarse);
-}
-
 void GridScalingKernelManager::runCoarseToFineKernelLB(const int level, InterpolationCells* coarseToFine, ICellNeigh &neighborFineToCoarse, CudaStreamIndex streamIndex) const
 {
     cudaStream_t stream = para->getStreamManager()->getStream(streamIndex);
     this->scalingCoarseToFine(para->getParD(level).get(), para->getParD(level+1).get(), coarseToFine, neighborFineToCoarse, stream);
 }
-
-void GridScalingKernelManager::runCoarseToFineKernelAD(const int level) const
-{
-    scaleCoarseToFineAdvectionDiffusion(
-        para->getParD(level)->distributions.f[0],
-        para->getParD(level+1)->distributions.f[0],
-        para->getParD(level)->distributionsAD.f[0],
-        para->getParD(level+1)->distributionsAD.f[0],
-        para->getParD(level)->neighborX,
-        para->getParD(level)->neighborY,
-        para->getParD(level)->neighborZ,
-        para->getParD(level+1)->neighborX,
-        para->getParD(level+1)->neighborY,
-        para->getParD(level+1)->neighborZ,
-        para->getParD(level)->numberOfNodes,
-        para->getParD(level+1)->numberOfNodes,
-        para->getParD(level)->isEvenTimestep,
-        para->getParD(level)->coarseToFine.coarseCellIndices,
-        para->getParD(level)->coarseToFine.fineCellIndices,
-        para->getParD(level)->coarseToFine.numberOfCells,
-        para->getParD(level)->viscosity,
-        para->getParD(level+1)->diffusivity,
-        para->getParD(level)->numberofthreads,
-        para->getParD(level)->neighborCoarseToFine);
-}

src/gpu/core/GridScaling/GridScalingKernelManager.h

@@ -62,15 +62,9 @@ public:
     //! \brief calls the device function of the fine to coarse grid interpolation kernelH
     void runFineToCoarseKernelLB(const int level, InterpolationCells *fineToCoarse, ICellNeigh &neighborFineToCoarse, CudaStreamIndex streamIndex) const;
 
-    //! \brief calls the device function of the fine to coarse grid interpolation kernel (advection diffusion)
-    void runFineToCoarseKernelAD(const int level) const;
-
     //! \brief calls the device function of the coarse to fine grid interpolation kernel
     void runCoarseToFineKernelLB(const int level, InterpolationCells *coarseToFine, ICellNeigh &neighborCoarseToFine, CudaStreamIndex streamIndex) const;
 
-    //! \brief calls the device function of the coarse to fine grid interpolation kernel (advection diffusion)
-    void runCoarseToFineKernelAD(const int level) const;
-
 private:
     //! \brief check if grid scaling was set
     //! \throws std::runtime_error if interpolation nodes were assigned, but no scaling function was set in the grid

src/gpu/core/GridScaling/Scaling.cu

@@ -47,6 +47,7 @@ using namespace vf::lbm::dir;
 using namespace vf::gpu;
 
 
+//////////////////////////////////////////////////////////////////////////
 // coarse to fine
 template <bool hasTurbulentViscosity>
 __global__ void scaleCoarseToFineCompressible_Device(
@@ -71,28 +72,7 @@ __global__ void scaleCoarseToFineCompressible_Device(
     ICellNeigh offsetCF);
 
 
-__global__ void scaleCoarseToFineAdvectionDiffusion_Device(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posCSWB,
-    uint* posFSWB,
-    uint kCF,
-    real nu,
-    real diffusivity_fine,
-    ICellNeigh neighborCoarseToFine);
-
-
+//////////////////////////////////////////////////////////////////////////
 // fine to coarse
 template <bool hasTurbulentViscosity>
 __global__ void scaleFineToCoarseCompressible_Device(
@@ -117,29 +97,8 @@ __global__ void scaleFineToCoarseCompressible_Device(
     ICellNeigh offsetFC);
 
 
-__global__ void scaleFineToCoarseAdvectionDiffusion_Device(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posC,
-    uint* posFSWB,
-    uint kFC,
-    real nu,
-    real diffusivity_coarse,
-    ICellNeigh neighborFineToCoarse);
-
 //////////////////////////////////////////////////////////////////////////
-
+// coarse to fine
 template <bool hasTurbulentViscosity>
 void scaleCoarseToFineCompressible(
     LBMSimulationParameter* parameterDeviceC,
@@ -187,54 +146,9 @@ template void scaleCoarseToFineCompressible<false>(
     ICellNeigh& neighborCoarseToFine,
     CUstream_st* stream);
 
-void scaleCoarseToFineAdvectionDiffusion(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posCSWB,
-    uint* posFSWB,
-    uint kCF,
-    real nu,
-    real diffusivity_fine,
-    uint numberOfThreads,
-    ICellNeigh neighborCoarseToFine)
-{
-    vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, kCF);
-
-    scaleCoarseToFineAdvectionDiffusion_Device<<<grid.grid, grid.threads>>>(
-        DC,
-        DF,
-        DD27C,
-        DD27F,
-        neighborCX,
-        neighborCY,
-        neighborCZ,
-        neighborFX,
-        neighborFY,
-        neighborFZ,
-        numberOfLBnodesC,
-        numberOfLBnodesF,
-        isEvenTimestep,
-        posCSWB,
-        posFSWB,
-        kCF,
-        nu,
-        diffusivity_fine,
-        neighborCoarseToFine);
-    getLastCudaError("scaleCoarseToFineAdvectionDiffusion_Device execution failed");
-}
-
 
+//////////////////////////////////////////////////////////////////////////
+// fine to coarse
 template <bool hasTurbulentViscosity>
 void scaleFineToCoarseCompressible(
     LBMSimulationParameter* parameterDeviceC,
@@ -281,51 +195,3 @@ template void scaleFineToCoarseCompressible<false>(
     ICells* fineToCoarse,
     ICellNeigh& neighborFineToCoarse,
     CUstream_st* stream);
-
-
-void scaleFineToCoarseAdvectionDiffusion(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posC,
-    uint* posFSWB,
-    uint kFC,
-    real nu,
-    real diffusivity_coarse,
-    uint numberOfThreads,
-    ICellNeigh neighborFineToCoarse)
-{
-    vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, kFC);
-
-    scaleFineToCoarseAdvectionDiffusion_Device<<<grid.grid, grid.threads>>>(
-        DC,
-        DF,
-        DD27C,
-        DD27F,
-        neighborCX,
-        neighborCY,
-        neighborCZ,
-        neighborFX,
-        neighborFY,
-        neighborFZ,
-        numberOfLBnodesC,
-        numberOfLBnodesF,
-        isEvenTimestep,
-        posC,
-        posFSWB,
-        kFC,
-        nu,
-        diffusivity_coarse,
-        neighborFineToCoarse);
-    getLastCudaError("scaleFineToCoarseAdvectionDiffusion_Device execution failed");
-}

src/gpu/core/GridScaling/Scaling.cuh

@@ -54,49 +54,4 @@ void scaleFineToCoarseCompressible(
     ICellNeigh& neighborFineToCoarse,
     CUstream_st* stream);
 
-
-void scaleCoarseToFineAdvectionDiffusion(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posCSWB,
-    uint* posFSWB,
-    uint kCF,
-    real nu,
-    real diffusivity_fine,
-    uint numberOfThreads,
-    ICellNeigh neighborCoarseToFine);
-
-void scaleFineToCoarseAdvectionDiffusion(
-    real* DC,
-    real* DF,
-    real* DD27C,
-    real* DD27F,
-    uint* neighborCX,
-    uint* neighborCY,
-    uint* neighborCZ,
-    uint* neighborFX,
-    uint* neighborFY,
-    uint* neighborFZ,
-    unsigned long long numberOfLBnodesC,
-    unsigned long long numberOfLBnodesF,
-    bool isEvenTimestep,
-    uint* posC,
-    uint* posFSWB,
-    uint kFC,
-    real nu,
-    real diffusivity_coarse,
-    uint numberOfThreads,
-    ICellNeigh neighborFineToCoarse);
-
 #endif