1//=======================================================================================
2// ____ ____ __ ______ __________ __ __ __ __
3// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
4// \ \ | | | | | |_) | | | | | | | / \ | |
5// \ \ | | | | | _ / | | | | | | / /\ \ | |
6// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
7// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
8// \ \ | | ________________________________________________________________
9// \ \ | | | ______________________________________________________________|
10// \ \| | | | __ __ __ __ ______ _______
11// \ | | |_____ | | | | | | | | | _ \ / _____)
12// \ | | _____| | | | | | | | | | | \ \ \_______
13// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
14// \ _____| |__| |________| \_______/ |__| |______/ (_______/
16// This file is part of VirtualFluids. VirtualFluids is free software: you can
17// redistribute it and/or modify it under the terms of the GNU General Public
18// License as published by the Free Software Foundation, either version 3 of
19// the License, or (at your option) any later version.
21// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
22// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
23// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
26// SPDX-License-Identifier: GPL-3.0-or-later
27// SPDX-FileCopyrightText: Copyright © VirtualFluids Project contributors, see AUTHORS.md in root folder
29//! \addtogroup gpu_BoundaryConditions BoundaryConditions
30//! \ingroup gpu_core core
32//! \author Martin Schoenherr
33//=======================================================================================
35#include <helper_cuda.h>
37#include "Calculation/Calculation.h"
38#include <cuda_helper/CudaGrid.h>
40#include "BoundaryConditions/BoundaryConditionFactory.h"
41#include "BoundaryConditions/Stress/Stress_Device.cuh"
42#include "Parameter/Parameter.h"
46using StressBC = BoundaryConditionFactory::StressBC;
48GridParameter getStressBCGridParameter(LBMSimulationParameter* parameterDevice)
50 return { parameterDevice->distributions.f[0], parameterDevice->omega, parameterDevice->turbulentViscosity,
51 parameterDevice->velocityX, parameterDevice->velocityY, parameterDevice->velocityZ,
52 parameterDevice->neighborX, parameterDevice->neighborY, parameterDevice->neighborZ,
53 parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep };
55template <StressBC stressBC, bool useDelayedBounceBack>
56void Stress(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
58 const vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
60 StressDevice27<stressBC, useDelayedBounceBack><<<grid.grid, grid.threads>>>(
61 getStressBCGridParameter(parameterDevice), *boundaryCondition, parameterDevice->momentumWallModel);
62 getLastCudaError("StressDevice27 execution failed");
65void StressBounceBackCompressible(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
67 Stress<StressBC::StressBounceBackCompressible, false>(parameterDevice, boundaryCondition);
70void StressBounceBackWithPressureCompressible(LBMSimulationParameter* parameterDevice,
71 QforBoundaryConditions* boundaryCondition)
73 Stress<StressBC::StressBounceBackWithPressureCompressible, false>(parameterDevice, boundaryCondition);
76void StressInterpolatedCompressible(LBMSimulationParameter* parameterDevice, QforBoundaryConditions* boundaryCondition)
78 Stress<StressBC::StressInterpolatedCompressible, false>(parameterDevice, boundaryCondition);