Skip to content
Snippets Groups Projects
Commit 3559b456 authored by Martin Schönherr's avatar Martin Schönherr :speech_balloon:
Browse files

move advection diffusion boundary conditions to BoundaryConditions folder and delete: 

QADPress7;
QADPress27;
QADPressNEQNeighbor27;
QADVel7;
QADVel27;
QAD7;
QNoSlipADIncomp7;
QNoSlipADIncomp27;
QADVeloIncomp7;
QADVeloIncomp27;
QADPressIncomp7;
QADPressIncomp27;
parent 8b182e65
No related branches found
No related tags found
1 merge request!350advection diffusion boundary conditions
Expand all Hide whitespace changes
Inline Side-by-side
Showing
with 1600 additions and 8248 deletions
src/gpu/core/BoundaryConditions/AdvectionDiffusion/AdvectionDiffusion.cu 0 → 100644
+ 148
− 0
View file @ 3559b456
//=======================================================================================
// ____ ____ __ ______ __________ __ __ __ __
// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
// \ \ | | | | | |_) | | | | | | | / \ | |
// \ \ | | | | | _ / | | | | | | / /\ \ | |
// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
// \ \ | | ________________________________________________________________
// \ \ | | | ______________________________________________________________|
// \ \| | | | __ __ __ __ ______ _______
// \ | | |_____ | | | | | | | | | _ \ / _____)
// \ | | _____| | | | | | | | | | | \ \ \_______
// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
// \ _____| |__| |________| \_______/ |__| |______/ (_______/
//
// This file is part of VirtualFluids. VirtualFluids is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \author Martin Schoenherr
//=======================================================================================
// includes, cuda
#include <cuda_runtime.h>
#include <helper_functions.h>
#include <helper_cuda.h>
#include "LBM/LB.h"
#include <cuda_helper/CudaGrid.h>
#include "BoundaryConditions/AdvectionDiffusion/AdvectionDiffusion_Device.cuh"
#include "Parameter/Parameter.h"
void AdvectionDiffusionSlipVelocityCompressible(
uint numberOfThreads,
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
AdvectionDiffusionSlipVelocityCompressible_Device <<< grid.grid, grid.threads >>> (
normalX,
normalY,
normalZ,
distributions,
distributionsAD,
QindexArray,
Qarrays,
numberOfBCnodes,
omegaDiffusivity,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("AdvectionDiffusionSlipVelocityCompressible_Device execution failed");
}
void AdvectionDiffusionDirichlet(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
AdvectionDiffusionDirichlet_Device<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("AdvectionDiffusionDirichlet_Device execution failed");
}
void AdvectionDiffusionBounceBack(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
AdvectionDiffusionBounceBack_Device<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("AdvectionDiffusionBounceBack_Device execution failed");
}
src/gpu/core/BoundaryConditions/AdvectionDiffusion/AdvectionDiffusion.h 0 → 100644
+ 93
− 0
View file @ 3559b456
//=======================================================================================
// ____ ____ __ ______ __________ __ __ __ __
// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
// \ \ | | | | | |_) | | | | | | | / \ | |
// \ \ | | | | | _ / | | | | | | / /\ \ | |
// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
// \ \ | | ________________________________________________________________
// \ \ | | | ______________________________________________________________|
// \ \| | | | __ __ __ __ ______ _______
// \ | | |_____ | | | | | | | | | _ \ / _____)
// \ | | _____| | | | | | | | | | | \ \ \_______
// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
// \ _____| |__| |________| \_______/ |__| |______/ (_______/
//
// This file is part of VirtualFluids. VirtualFluids is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \author Martin Schoenherr
//=======================================================================================
#ifndef AdvectionDiffusion_H
#define AdvectionDiffusion_H
#include "LBM/LB.h"
#include <cuda.h>
#include <cuda_runtime.h>
struct LBMSimulationParameter;
class Parameter;
//////////////////////////////////////////////////////////////////////////
//! \brief defines the behavior of a slip-AD boundary condition
void AdvectionDiffusionSlipVelocityCompressible(
uint numberOfThreads,
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void AdvectionDiffusionDirichlet(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void AdvectionDiffusionBounceBack(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
#endif
src/gpu/core/BoundaryConditions/AdvectionDiffusion/AdvectionDiffusionBounceBack.cu 0 → 100644
+ 306
− 0
View file @ 3559b456
//=======================================================================================
// ____ ____ __ ______ __________ __ __ __ __
// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
// \ \ | | | | | |_) | | | | | | | / \ | |
// \ \ | | | | | _ / | | | | | | / /\ \ | |
// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
// \ \ | | ________________________________________________________________
// \ \ | | | ______________________________________________________________|
// \ \| | | | __ __ __ __ ______ _______
// \ | | |_____ | | | | | | | | | _ \ / _____)
// \ | | _____| | | | | | | | | | | \ \ \_______
// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
// \ _____| |__| |________| \_______/ |__| |______/ (_______/
//
// This file is part of VirtualFluids. VirtualFluids is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \author Martin Schoenherr
//=======================================================================================
#include "LBM/LB.h"
#include "lbm/constants/D3Q27.h"
#include <basics/constants/NumericConstants.h>
using namespace vf::basics::constant;
using namespace vf::lbm::dir;
__global__ void AdvectionDiffusionBounceBack_Device(
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
Distributions27 D27;
if (isEvenTimestep==true)
{
D27.f[dP00] = &DD27[dP00 * numberOfLBnodes];
D27.f[dM00] = &DD27[dM00 * numberOfLBnodes];
D27.f[d0P0] = &DD27[d0P0 * numberOfLBnodes];
D27.f[d0M0] = &DD27[d0M0 * numberOfLBnodes];
D27.f[d00P] = &DD27[d00P * numberOfLBnodes];
D27.f[d00M] = &DD27[d00M * numberOfLBnodes];
D27.f[dPP0] = &DD27[dPP0 * numberOfLBnodes];
D27.f[dMM0] = &DD27[dMM0 * numberOfLBnodes];
D27.f[dPM0] = &DD27[dPM0 * numberOfLBnodes];
D27.f[dMP0] = &DD27[dMP0 * numberOfLBnodes];
D27.f[dP0P] = &DD27[dP0P * numberOfLBnodes];
D27.f[dM0M] = &DD27[dM0M * numberOfLBnodes];
D27.f[dP0M] = &DD27[dP0M * numberOfLBnodes];
D27.f[dM0P] = &DD27[dM0P * numberOfLBnodes];
D27.f[d0PP] = &DD27[d0PP * numberOfLBnodes];
D27.f[d0MM] = &DD27[d0MM * numberOfLBnodes];
D27.f[d0PM] = &DD27[d0PM * numberOfLBnodes];
D27.f[d0MP] = &DD27[d0MP * numberOfLBnodes];
D27.f[d000] = &DD27[d000 * numberOfLBnodes];
D27.f[dPPP] = &DD27[dPPP * numberOfLBnodes];
D27.f[dMMP] = &DD27[dMMP * numberOfLBnodes];
D27.f[dPMP] = &DD27[dPMP * numberOfLBnodes];
D27.f[dMPP] = &DD27[dMPP * numberOfLBnodes];
D27.f[dPPM] = &DD27[dPPM * numberOfLBnodes];
D27.f[dMMM] = &DD27[dMMM * numberOfLBnodes];
D27.f[dPMM] = &DD27[dPMM * numberOfLBnodes];
D27.f[dMPM] = &DD27[dMPM * numberOfLBnodes];
}
else
{
D27.f[dM00] = &DD27[dP00 * numberOfLBnodes];
D27.f[dP00] = &DD27[dM00 * numberOfLBnodes];
D27.f[d0M0] = &DD27[d0P0 * numberOfLBnodes];
D27.f[d0P0] = &DD27[d0M0 * numberOfLBnodes];
D27.f[d00M] = &DD27[d00P * numberOfLBnodes];
D27.f[d00P] = &DD27[d00M * numberOfLBnodes];
D27.f[dMM0] = &DD27[dPP0 * numberOfLBnodes];
D27.f[dPP0] = &DD27[dMM0 * numberOfLBnodes];
D27.f[dMP0] = &DD27[dPM0 * numberOfLBnodes];
D27.f[dPM0] = &DD27[dMP0 * numberOfLBnodes];
D27.f[dM0M] = &DD27[dP0P * numberOfLBnodes];
D27.f[dP0P] = &DD27[dM0M * numberOfLBnodes];
D27.f[dM0P] = &DD27[dP0M * numberOfLBnodes];
D27.f[dP0M] = &DD27[dM0P * numberOfLBnodes];
D27.f[d0MM] = &DD27[d0PP * numberOfLBnodes];
D27.f[d0PP] = &DD27[d0MM * numberOfLBnodes];
D27.f[d0MP] = &DD27[d0PM * numberOfLBnodes];
D27.f[d0PM] = &DD27[d0MP * numberOfLBnodes];
D27.f[d000] = &DD27[d000 * numberOfLBnodes];
D27.f[dPPP] = &DD27[dMMM * numberOfLBnodes];
D27.f[dMMP] = &DD27[dPPM * numberOfLBnodes];
D27.f[dPMP] = &DD27[dMPM * numberOfLBnodes];
D27.f[dMPP] = &DD27[dPMM * numberOfLBnodes];
D27.f[dPPM] = &DD27[dMMP * numberOfLBnodes];
D27.f[dMMM] = &DD27[dPPP * numberOfLBnodes];
D27.f[dPMM] = &DD27[dMPP * numberOfLBnodes];
D27.f[dMPM] = &DD27[dPMP * numberOfLBnodes];
}
////////////////////////////////////////////////////////////////////////////////
const unsigned x = threadIdx.x;
const unsigned y = blockIdx.x;
const unsigned z = blockIdx.y;
const unsigned nx = blockDim.x;
const unsigned ny = gridDim.x;
const unsigned k = nx*(ny*z + y) + x;
//////////////////////////////////////////////////////////////////////////
if(k<numberOfBCnodes)
{
////////////////////////////////////////////////////////////////////////////////
real *q_dirE, *q_dirW, *q_dirN, *q_dirS, *q_dirT, *q_dirB,
*q_dirNE, *q_dirSW, *q_dirSE, *q_dirNW, *q_dirTE, *q_dirBW,
*q_dirBE, *q_dirTW, *q_dirTN, *q_dirBS, *q_dirBN, *q_dirTS,
*q_dirTNE, *q_dirTSW, *q_dirTSE, *q_dirTNW, *q_dirBNE, *q_dirBSW,
*q_dirBSE, *q_dirBNW;
q_dirE = &QQ[dP00 * numberOfBCnodes];
q_dirW = &QQ[dM00 * numberOfBCnodes];
q_dirN = &QQ[d0P0 * numberOfBCnodes];
q_dirS = &QQ[d0M0 * numberOfBCnodes];
q_dirT = &QQ[d00P * numberOfBCnodes];
q_dirB = &QQ[d00M * numberOfBCnodes];
q_dirNE = &QQ[dPP0 * numberOfBCnodes];
q_dirSW = &QQ[dMM0 * numberOfBCnodes];
q_dirSE = &QQ[dPM0 * numberOfBCnodes];
q_dirNW = &QQ[dMP0 * numberOfBCnodes];
q_dirTE = &QQ[dP0P * numberOfBCnodes];
q_dirBW = &QQ[dM0M * numberOfBCnodes];
q_dirBE = &QQ[dP0M * numberOfBCnodes];
q_dirTW = &QQ[dM0P * numberOfBCnodes];
q_dirTN = &QQ[d0PP * numberOfBCnodes];
q_dirBS = &QQ[d0MM * numberOfBCnodes];
q_dirBN = &QQ[d0PM * numberOfBCnodes];
q_dirTS = &QQ[d0MP * numberOfBCnodes];
q_dirTNE = &QQ[dPPP * numberOfBCnodes];
q_dirTSW = &QQ[dMMP * numberOfBCnodes];
q_dirTSE = &QQ[dPMP * numberOfBCnodes];
q_dirTNW = &QQ[dMPP * numberOfBCnodes];
q_dirBNE = &QQ[dPPM * numberOfBCnodes];
q_dirBSW = &QQ[dMMM * numberOfBCnodes];
q_dirBSE = &QQ[dPMM * numberOfBCnodes];
q_dirBNW = &QQ[dMPM * numberOfBCnodes];
////////////////////////////////////////////////////////////////////////////////
//index
unsigned int KQK = k_Q[k];
unsigned int ke = KQK;
unsigned int kw = neighborX[KQK];
unsigned int kn = KQK;
unsigned int ks = neighborY[KQK];
unsigned int kt = KQK;
unsigned int kb = neighborZ[KQK];
unsigned int ksw = neighborY[kw];
unsigned int kne = KQK;
unsigned int kse = ks;
unsigned int knw = kw;
unsigned int kbw = neighborZ[kw];
unsigned int kte = KQK;
unsigned int kbe = kb;
unsigned int ktw = kw;
unsigned int kbs = neighborZ[ks];
unsigned int ktn = KQK;
unsigned int kbn = kb;
unsigned int kts = ks;
unsigned int ktse = ks;
unsigned int kbnw = kbw;
unsigned int ktnw = kw;
unsigned int kbse = kbs;
unsigned int ktsw = ksw;
unsigned int kbne = kb;
unsigned int ktne = KQK;
unsigned int kbsw = neighborZ[ksw];
////////////////////////////////////////////////////////////////////////////////
real f27_W = (D27.f[dP00])[ke ];
real f27_E = (D27.f[dM00])[kw ];
real f27_S = (D27.f[d0P0])[kn ];
real f27_N = (D27.f[d0M0])[ks ];
real f27_B = (D27.f[d00P])[kt ];
real f27_T = (D27.f[d00M])[kb ];
real f27_SW = (D27.f[dPP0])[kne ];
real f27_NE = (D27.f[dMM0])[ksw ];
real f27_NW = (D27.f[dPM0])[kse ];
real f27_SE = (D27.f[dMP0])[knw ];
real f27_BW = (D27.f[dP0P])[kte ];
real f27_TE = (D27.f[dM0M])[kbw ];
real f27_TW = (D27.f[dP0M])[kbe ];
real f27_BE = (D27.f[dM0P])[ktw ];
real f27_BS = (D27.f[d0PP])[ktn ];
real f27_TN = (D27.f[d0MM])[kbs ];
real f27_TS = (D27.f[d0PM])[kbn ];
real f27_BN = (D27.f[d0MP])[kts ];
real f27_BSW = (D27.f[dPPP])[ktne ];
real f27_BNE = (D27.f[dMMP])[ktsw ];
real f27_BNW = (D27.f[dPMP])[ktse ];
real f27_BSE = (D27.f[dMPP])[ktnw ];
real f27_TSW = (D27.f[dPPM])[kbne ];
real f27_TNE = (D27.f[dMMM])[kbsw ];
real f27_TNW = (D27.f[dPMM])[kbse ];
real f27_TSE = (D27.f[dMPM])[kbnw ];
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
if (isEvenTimestep==false)
{
D27.f[dP00] = &DD27[dP00 * numberOfLBnodes];
D27.f[dM00] = &DD27[dM00 * numberOfLBnodes];
D27.f[d0P0] = &DD27[d0P0 * numberOfLBnodes];
D27.f[d0M0] = &DD27[d0M0 * numberOfLBnodes];
D27.f[d00P] = &DD27[d00P * numberOfLBnodes];
D27.f[d00M] = &DD27[d00M * numberOfLBnodes];
D27.f[dPP0] = &DD27[dPP0 * numberOfLBnodes];
D27.f[dMM0] = &DD27[dMM0 * numberOfLBnodes];
D27.f[dPM0] = &DD27[dPM0 * numberOfLBnodes];
D27.f[dMP0] = &DD27[dMP0 * numberOfLBnodes];
D27.f[dP0P] = &DD27[dP0P * numberOfLBnodes];
D27.f[dM0M] = &DD27[dM0M * numberOfLBnodes];
D27.f[dP0M] = &DD27[dP0M * numberOfLBnodes];
D27.f[dM0P] = &DD27[dM0P * numberOfLBnodes];
D27.f[d0PP] = &DD27[d0PP * numberOfLBnodes];
D27.f[d0MM] = &DD27[d0MM * numberOfLBnodes];
D27.f[d0PM] = &DD27[d0PM * numberOfLBnodes];
D27.f[d0MP] = &DD27[d0MP * numberOfLBnodes];
D27.f[d000] = &DD27[d000 * numberOfLBnodes];
D27.f[dPPP] = &DD27[dPPP * numberOfLBnodes];
D27.f[dMMP] = &DD27[dMMP * numberOfLBnodes];
D27.f[dPMP] = &DD27[dPMP * numberOfLBnodes];
D27.f[dMPP] = &DD27[dMPP * numberOfLBnodes];
D27.f[dPPM] = &DD27[dPPM * numberOfLBnodes];
D27.f[dMMM] = &DD27[dMMM * numberOfLBnodes];
D27.f[dPMM] = &DD27[dPMM * numberOfLBnodes];
D27.f[dMPM] = &DD27[dMPM * numberOfLBnodes];
}
else
{
D27.f[dM00] = &DD27[dP00 * numberOfLBnodes];
D27.f[dP00] = &DD27[dM00 * numberOfLBnodes];
D27.f[d0M0] = &DD27[d0P0 * numberOfLBnodes];
D27.f[d0P0] = &DD27[d0M0 * numberOfLBnodes];
D27.f[d00M] = &DD27[d00P * numberOfLBnodes];
D27.f[d00P] = &DD27[d00M * numberOfLBnodes];
D27.f[dMM0] = &DD27[dPP0 * numberOfLBnodes];
D27.f[dPP0] = &DD27[dMM0 * numberOfLBnodes];
D27.f[dMP0] = &DD27[dPM0 * numberOfLBnodes];
D27.f[dPM0] = &DD27[dMP0 * numberOfLBnodes];
D27.f[dM0M] = &DD27[dP0P * numberOfLBnodes];
D27.f[dP0P] = &DD27[dM0M * numberOfLBnodes];
D27.f[dM0P] = &DD27[dP0M * numberOfLBnodes];
D27.f[dP0M] = &DD27[dM0P * numberOfLBnodes];
D27.f[d0MM] = &DD27[d0PP * numberOfLBnodes];
D27.f[d0PP] = &DD27[d0MM * numberOfLBnodes];
D27.f[d0MP] = &DD27[d0PM * numberOfLBnodes];
D27.f[d0PM] = &DD27[d0MP * numberOfLBnodes];
D27.f[d000] = &DD27[d000 * numberOfLBnodes];
D27.f[dPPP] = &DD27[dMMM * numberOfLBnodes];
D27.f[dMMP] = &DD27[dPPM * numberOfLBnodes];
D27.f[dPMP] = &DD27[dMPM * numberOfLBnodes];
D27.f[dMPP] = &DD27[dPMM * numberOfLBnodes];
D27.f[dPPM] = &DD27[dMMP * numberOfLBnodes];
D27.f[dMMM] = &DD27[dPPP * numberOfLBnodes];
D27.f[dPMM] = &DD27[dMPP * numberOfLBnodes];
D27.f[dMPM] = &DD27[dPMP * numberOfLBnodes];
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
real q;
q = q_dirE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dM00])[kw ]=f27_E ;
q = q_dirW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dP00])[ke ]=f27_W ;
q = q_dirN[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0M0])[ks ]=f27_N ;
q = q_dirS[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0P0])[kn ]=f27_S ;
q = q_dirT[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d00M])[kb ]=f27_T ;
q = q_dirB[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d00P])[kt ]=f27_B ;
q = q_dirNE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMM0])[ksw ]=f27_NE ;
q = q_dirSW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPP0])[kne ]=f27_SW ;
q = q_dirSE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMP0])[knw ]=f27_SE ;
q = q_dirNW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPM0])[kse ]=f27_NW ;
q = q_dirTE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dM0M])[kbw ]=f27_TE ;
q = q_dirBW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dP0P])[kte ]=f27_BW ;
q = q_dirBE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dM0P])[ktw ]=f27_BE ;
q = q_dirTW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dP0M])[kbe ]=f27_TW ;
q = q_dirTN[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0MM])[kbs ]=f27_TN ;
q = q_dirBS[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0PP])[ktn ]=f27_BS ;
q = q_dirBN[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0MP])[kts ]=f27_BN ;
q = q_dirTS[k]; if (q>=c0o1 && q<=c1o1) (D27.f[d0PM])[kbn ]=f27_TS ;
q = q_dirTNE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMMM])[kbsw]=f27_TNE;
q = q_dirBSW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPPP])[ktne]=f27_BSW;
q = q_dirBNE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMMP])[ktsw]=f27_BNE;
q = q_dirTSW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPPM])[kbne]=f27_TSW;
q = q_dirTSE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMPM])[kbnw]=f27_TSE;
q = q_dirBNW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPMP])[ktse]=f27_BNW;
q = q_dirBSE[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dMPP])[ktnw]=f27_BSE;
q = q_dirTNW[k]; if (q>=c0o1 && q<=c1o1) (D27.f[dPMM])[kbse]=f27_TNW;
}
}
src/gpu/core/BoundaryConditions/AdvectionDiffusion/AdvectionDiffusion_Device.cuh 0 → 100644
+ 84
− 0
View file @ 3559b456
//=======================================================================================
// ____ ____ __ ______ __________ __ __ __ __
// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
// \ \ | | | | | |_) | | | | | | | / \ | |
// \ \ | | | | | _ / | | | | | | / /\ \ | |
// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
// \ \ | | ________________________________________________________________
// \ \ | | | ______________________________________________________________|
// \ \| | | | __ __ __ __ ______ _______
// \ | | |_____ | | | | | | | | | _ \ / _____)
// \ | | _____| | | | | | | | | | | \ \ \_______
// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
// \ _____| |__| |________| \_______/ |__| |______/ (_______/
//
// This file is part of VirtualFluids. VirtualFluids is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \author Martin Schoenherr
//=======================================================================================
#ifndef AdvectionDiffusion_Device_H
#define AdvectionDiffusion_Device_H
#include "LBM/LB.h"
//////////////////////////////////////////////////////////////////////////
//! \brief \ref AD_SlipVelDeviceComp : device function for the slip-AD boundary condition
__global__ void AdvectionDiffusionSlipVelocityCompressible_Device(
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void AdvectionDiffusionDirichlet_Device(
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void AdvectionDiffusionBounceBack_Device(
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
#endif
src/gpu/core/Calculation/UpdateGrid27.cpp
+ 31
− 9
View file @ 3559b456
//=======================================================================================
// ____ ____ __ ______ __________ __ __ __ __
// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
// \ \ | | | | | |_) | | | | | | | / \ | |
// \ \ | | | | | _ / | | | | | | / /\ \ | |
// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
// \ \ | | ________________________________________________________________
// \ \ | | | ______________________________________________________________|
// \ \| | | | __ __ __ __ ______ _______
// \ | | |_____ | | | | | | | | | _ \ / _____)
// \ | | _____| | | | | | | | | | | \ \ \_______
// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
// \ _____| |__| |________| \_______/ |__| |______/ (_______/
//
// This file is part of VirtualFluids. VirtualFluids is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \author Martin Schoenherr, Stephan Lenz, Anna Wellmann
//=======================================================================================
#include "UpdateGrid27.h" #include "UpdateGrid27.h"
#include <cuda_runtime.h> #include <cuda_runtime.h>
...@@ -127,13 +157,6 @@ void UpdateGrid27::exchangeMultiGPU(int level, CudaStreamIndex streamIndex) ...@@ -127,13 +157,6 @@ void UpdateGrid27::exchangeMultiGPU(int level, CudaStreamIndex streamIndex)
exchangePostCollDataADZGPU27(para.get(), comm, cudaMemoryManager.get(), level); exchangePostCollDataADZGPU27(para.get(), comm, cudaMemoryManager.get(), level);
} }
//////////////////////////////////////////////////////////////////////////
// D E P R E C A T E D
//////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////
// 1D domain decomposition
// exchangePostCollDataGPU27(para, comm, level);
} }
void UpdateGrid27::exchangeMultiGPU_noStreams_withPrepare(int level, bool useReducedComm) void UpdateGrid27::exchangeMultiGPU_noStreams_withPrepare(int level, bool useReducedComm)
{ {
...@@ -236,9 +259,8 @@ void UpdateGrid27::postCollisionBC(int level, uint t) ...@@ -236,9 +259,8 @@ void UpdateGrid27::postCollisionBC(int level, uint t)
if (para->getDiffOn()) if (para->getDiffOn())
{ {
this->adKernelManager->runADgeometryBCKernel(level); this->adKernelManager->runADgeometryBCKernel(level);
this->adKernelManager->runADveloBCKernel(level); this->adKernelManager->runADDirichletBCKernel(level);
this->adKernelManager->runADslipBCKernel(level); this->adKernelManager->runADslipBCKernel(level);
this->adKernelManager->runADpressureBCKernel(level);
} }
} }
......
src/gpu/core/GPU/AdvectionDiffusionBCs27.cu deleted 100644 → 0
+ 0
− 7147
View file @ 8b182e65
This diff is collapsed.
src/gpu/core/GPU/GPU_Interface.h
+ 0
− 238
View file @ 3559b456
...@@ -322,244 +322,6 @@ void LBCalcMeasurePoints27(real* vxMP, ...@@ -322,244 +322,6 @@ void LBCalcMeasurePoints27(real* vxMP,
unsigned int numberOfThreads, unsigned int numberOfThreads,
bool isEvenTimestep); bool isEvenTimestep);
void QADDev7(unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
//////////////////////////////////////////////////////////////////////////
//! \brief defines the behavior of a slip-AD boundary condition
void ADSlipVelDevComp(
uint numberOfThreads,
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADDirichletDev27( unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADBBDev27( unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADVelDev7(unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADVelDev27( unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADPressDev7( unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADPressDev27(unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADPressNEQNeighborDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
int* k_Q,
int* k_N,
int numberOfBCnodes,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep
);
void QNoSlipADincompDev7(unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QNoSlipADincompDev27(unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADVeloIncompDev7( unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADVeloIncompDev27( unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADPressIncompDev7( unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void QADPressIncompDev27( unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
void ScaleCF27( real* DC, void ScaleCF27( real* DC,
real* DF, real* DF,
unsigned int* neighborCX, unsigned int* neighborCX,
......
src/gpu/core/GPU/GPU_Kernels.cuh
+ 0
− 219
View file @ 3559b456
...@@ -293,227 +293,8 @@ __global__ void LBCalcMeasurePoints(real* vxMP, ...@@ -293,227 +293,8 @@ __global__ void LBCalcMeasurePoints(real* vxMP,
real* DD, real* DD,
bool isEvenTimestep); bool isEvenTimestep);
//Advection / Diffusion BCs
__global__ void QAD7( real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////
//! \brief \ref AD_SlipVelDeviceComp : device function for the slip-AD boundary condition
__global__ void AD_SlipVelDeviceComp(
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADDirichlet27( real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADBB27( real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADVel7( real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADVel27(real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADPress7( real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADPress27( real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADPressNEQNeighbor27(
real* DD,
real* DD27,
int* k_Q,
int* k_N,
int numberOfBCnodes,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep
);
__global__ void QNoSlipADincomp7( real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QNoSlipADincomp27( real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADVeloIncomp7( real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADVeloIncomp27( real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADPressIncomp7(real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
__global__ void QADPressIncomp27( real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
......
src/gpu/core/GPU/LBMKernel.cu
+ 0
− 541
View file @ 3559b456
...@@ -698,547 +698,6 @@ void LBCalcMeasurePoints27( ...@@ -698,547 +698,6 @@ void LBCalcMeasurePoints27(
getLastCudaError("LBCalcMeasurePoints execution failed"); getLastCudaError("LBCalcMeasurePoints execution failed");
} }
////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////
void QADPressDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADPress7<<< grid.grid, grid.threads >>>(
DD,
DD7,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADPress7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADPressDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADPress27<<< grid.grid, grid.threads >>>(
DD,
DD27,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADPress27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADPressNEQNeighborDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
int* k_Q,
int* k_N,
int numberOfBCnodes,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADPressNEQNeighbor27<<< grid.grid, grid.threads >>>(
DD,
DD27,
k_Q,
k_N,
numberOfBCnodes,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADPressNEQNeighbor27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADVelDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADVel7<<< grid.grid, grid.threads >>> (
DD,
DD7,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADVel7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADVelDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADVel27<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADVel27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QAD7<<< grid.grid, grid.threads >>> (
DD,
DD7,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QAD7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////
void ADSlipVelDevComp(
uint numberOfThreads,
real * normalX,
real * normalY,
real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfBCnodes,
real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
AD_SlipVelDeviceComp <<< grid.grid, grid.threads >>> (
normalX,
normalY,
normalZ,
distributions,
distributionsAD,
QindexArray,
Qarrays,
numberOfBCnodes,
omegaDiffusivity,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("AD_SlipVelDeviceComp execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADDirichletDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADDirichlet27<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADDirichletDev27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADBBDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADBB27<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADBB27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QNoSlipADincompDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QNoSlipADincomp7<<< grid.grid, grid.threads >>> (
DD,
DD7,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QNoSlipADincomp7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QNoSlipADincompDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QNoSlipADincomp27<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QNoSlipADincomp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADVeloIncompDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADVeloIncomp7<<< grid.grid, grid.threads >>> (
DD,
DD7,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADVeloIncomp7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADVeloIncompDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADVeloIncomp27<<< grid.grid, grid.threads >>> (
DD,
DD27,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADVeloIncomp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADPressIncompDev7(
unsigned int numberOfThreads,
real* DD,
real* DD7,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADPressIncomp7<<< grid.grid, grid.threads >>>(
DD,
DD7,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADPressIncomp7 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void QADPressIncompDev27(
unsigned int numberOfThreads,
real* DD,
real* DD27,
real* temp,
real* velo,
real diffusivity,
int* k_Q,
real* QQ,
unsigned int numberOfBCnodes,
real om1,
unsigned int* neighborX,
unsigned int* neighborY,
unsigned int* neighborZ,
unsigned long long numberOfLBnodes,
bool isEvenTimestep)
{
vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(numberOfThreads, numberOfBCnodes);
QADPressIncomp27<<< grid.grid, grid.threads >>>(
DD,
DD27,
temp,
velo,
diffusivity,
k_Q,
QQ,
numberOfBCnodes,
om1,
neighborX,
neighborY,
neighborZ,
numberOfLBnodes,
isEvenTimestep);
getLastCudaError("QADPressIncomp27 execution failed");
}
//////////////////////////////////////////////////////////////////////////
void ScaleCF27( void ScaleCF27(
real* DC, real* DC,
real* DF, real* DF,
......
src/gpu/core/Kernel/ADKernelManager.cpp
+ 5
− 88
View file @ 3559b456
...@@ -26,13 +26,12 @@ ...@@ -26,13 +26,12 @@
// You should have received a copy of the GNU General Public License along // You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. // with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
// //
//! \file ADKernelManager.h
//! \ingroup KernelManager
//! \author Martin Schoenherr //! \author Martin Schoenherr
//======================================================================================= //=======================================================================================
#include "Kernel/ADKernelManager.h" #include "Kernel/ADKernelManager.h"
#include "GPU/CudaMemoryManager.h" #include "GPU/CudaMemoryManager.h"
#include "GPU/GPU_Interface.h" #include "GPU/GPU_Interface.h"
#include "BoundaryConditions/AdvectionDiffusion/AdvectionDiffusion.h"
#include "Parameter/Parameter.h" #include "Parameter/Parameter.h"
#include "Kernel/AdvectionDiffusionKernel.h" #include "Kernel/AdvectionDiffusionKernel.h"
...@@ -69,7 +68,7 @@ void ADKernelManager::runADcollisionKernel(const int level)const ...@@ -69,7 +68,7 @@ void ADKernelManager::runADcollisionKernel(const int level)const
void ADKernelManager::runADslipBCKernel(const int level) const{ void ADKernelManager::runADslipBCKernel(const int level) const{
if (para->getParD(level)->slipBC.numberOfBCnodes > 1) { if (para->getParD(level)->slipBC.numberOfBCnodes > 1) {
ADSlipVelDevComp( AdvectionDiffusionSlipVelocityCompressible(
para->getParD(level)->numberofthreads, para->getParD(level)->numberofthreads,
para->getParD(level)->slipBC.normalX, para->getParD(level)->slipBC.normalX,
para->getParD(level)->slipBC.normalY, para->getParD(level)->slipBC.normalY,
...@@ -88,71 +87,10 @@ void ADKernelManager::runADslipBCKernel(const int level) const{ ...@@ -88,71 +87,10 @@ void ADKernelManager::runADslipBCKernel(const int level) const{
} }
} }
void ADKernelManager::runADpressureBCKernel(const int level) const{
if (para->getParD(level)->TempPress.kTemp > 0){
// QADPressIncompDev27(
// para->getParD(level)->numberofthreads,
// para->getParD(level)->distributions.f[0],
// para->getParD(level)->distributionsAD27.f[0],
// para->getParD(level)->TempPress.temp,
// para->getParD(level)->TempPress.velo,
// para->getParD(level)->diffusivity,
// para->getParD(level)->TempPress.k,
// para->getParD(level)->pressureBC.q27[0],
// para->getParD(level)->TempPress.kTemp,
// para->getParD(level)->omega,
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
// para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
//////////////////////////////////////////////////////////////////////////
// C O M P R E S S I B L E
//////////////////////////////////////////////////////////////////////////
QADPressDev27(
para->getParD(level)->numberofthreads,
para->getParD(level)->distributions.f[0],
para->getParD(level)->distributionsAD.f[0],
para->getParD(level)->TempPress.temp,
para->getParD(level)->TempPress.velo,
para->getParD(level)->diffusivity,
para->getParD(level)->TempPress.k,
para->getParD(level)->pressureBC.q27[0],
para->getParD(level)->TempPress.kTemp,
para->getParD(level)->omega,
para->getParD(level)->neighborX,
para->getParD(level)->neighborY,
para->getParD(level)->neighborZ,
para->getParD(level)->numberOfNodes,
para->getParD(level)->isEvenTimestep);
}
}
void ADKernelManager::runADgeometryBCKernel(const int level) const void ADKernelManager::runADgeometryBCKernel(const int level) const
{ {
if (para->getParD(level)->geometryBC.numberOfBCnodes > 0) { if (para->getParD(level)->geometryBC.numberOfBCnodes > 0) {
// QNoSlipADincompDev27( AdvectionDiffusionBounceBack(
// para->getParD(level)->numberofthreads,
// para->getParD(level)->distributions.f[0],
// para->getParD(level)->distributionsAD27.f[0],
// para->getParD(level)->Temp.temp,
// para->getParD(level)->diffusivity,
// para->getParD(level)->Temp.k,
// para->getParD(level)->geometryBC.q27[0],
// para->getParD(level)->Temp.kTemp,
// para->getParD(level)->omega,
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
// para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
//////////////////////////////////////////////////////////////////////////
// C O M P R E S S I B L E
//////////////////////////////////////////////////////////////////////////
QADBBDev27(
para->getParD(level)->numberofthreads, para->getParD(level)->numberofthreads,
para->getParD(level)->distributions.f[0], para->getParD(level)->distributions.f[0],
para->getParD(level)->distributionsAD.f[0], para->getParD(level)->distributionsAD.f[0],
...@@ -170,34 +108,13 @@ void ADKernelManager::runADgeometryBCKernel(const int level) const ...@@ -170,34 +108,13 @@ void ADKernelManager::runADgeometryBCKernel(const int level) const
} }
} }
void ADKernelManager::runADveloBCKernel(const int level) const{ void ADKernelManager::runADDirichletBCKernel(const int level) const{
if (para->getParD(level)->TempVel.kTemp > 0){ if (para->getParD(level)->TempVel.kTemp > 0){
// QADVeloIncompDev27( AdvectionDiffusionDirichlet(
// para->getParD(level)->numberofthreads,
// para->getParD(level)->distributions.f[0],
// para->getParD(level)->distributionsAD27.f[0],
// para->getParD(level)->TempVel.temp,
// para->getParD(level)->TempVel.velo,
// para->getParD(level)->diffusivity,
// para->getParD(level)->TempVel.k,
// para->getParD(level)->velocityBC.q27[0],
// para->getParD(level)->TempVel.kTemp,
// para->getParD(level)->omega,
// para->getParD(level)->neighborX,
// para->getParD(level)->neighborY,
// para->getParD(level)->neighborZ,
// para->getParD(level)->numberOfNodes,
// para->getParD(level)->isEvenTimestep);
//////////////////////////////////////////////////////////////////////////
// C O M P R E S S I B L E
//////////////////////////////////////////////////////////////////////////
QADVelDev27(
para->getParD(level)->numberofthreads, para->getParD(level)->numberofthreads,
para->getParD(level)->distributions.f[0], para->getParD(level)->distributions.f[0],
para->getParD(level)->distributionsAD.f[0], para->getParD(level)->distributionsAD.f[0],
para->getParD(level)->TempVel.tempPulse, para->getParD(level)->TempVel.tempPulse,
para->getParD(level)->TempVel.velo,
para->getParD(level)->diffusivity, para->getParD(level)->diffusivity,
para->getParD(level)->velocityBC.k, para->getParD(level)->velocityBC.k,
para->getParD(level)->velocityBC.q27[0], para->getParD(level)->velocityBC.q27[0],
......
src/gpu/core/Kernel/ADKernelManager.h
+ 1
− 6
View file @ 3559b456
...@@ -26,8 +26,6 @@ ...@@ -26,8 +26,6 @@
// You should have received a copy of the GNU General Public License along // You should have received a copy of the GNU General Public License along
// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. // with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
// //
//! \file ADKernelManager.h
//! \ingroup KernelManager
//! \author Martin Schoenherr //! \author Martin Schoenherr
//======================================================================================= //=======================================================================================
#ifndef ADVECTION_DIFFUSION_H #ifndef ADVECTION_DIFFUSION_H
...@@ -65,14 +63,11 @@ public: ...@@ -65,14 +63,11 @@ public:
void runADgeometryBCKernel(const int level) const; void runADgeometryBCKernel(const int level) const;
//! \brief calls the device function of the velocity boundary condition for advection diffusion //! \brief calls the device function of the velocity boundary condition for advection diffusion
void runADveloBCKernel(const int level) const; void runADDirichletBCKernel(const int level) const;
//! \brief calls the device function of the slip boundary condition for advection diffusion //! \brief calls the device function of the slip boundary condition for advection diffusion
void runADslipBCKernel(const int level) const; void runADslipBCKernel(const int level) const;
//! \brief calls the device function of the pressure boundary condition for advection diffusion
void runADpressureBCKernel(const int level) const;
//! \brief copy the concentration from device to host and writes VTK file with concentration //! \brief copy the concentration from device to host and writes VTK file with concentration
//! \param cudaMemoryManager instance of class CudaMemoryManager //! \param cudaMemoryManager instance of class CudaMemoryManager
void printAD(const int level, SPtr<CudaMemoryManager> cudaMemoryManager) const; void printAD(const int level, SPtr<CudaMemoryManager> cudaMemoryManager) const;
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment