From 9f659e73d83c2004f275de6544483bf5cd295c48 Mon Sep 17 00:00:00 2001
From: Martin Schoenherr <m.schoenherr@tu-braunschweig.de>
Date: Mon, 4 Dec 2023 10:53:46 +0100
Subject: [PATCH] delete boundary conditions:
PrecursorEquilibrium
PrecursorDistributionsSubgridDistance
---
.../BoundaryConditions/Precursor/Precursor.cu | 73 -----
.../BoundaryConditions/Precursor/Precursor.h | 4 -
.../PrecursorDistributionsSubgridDistance.cu | 212 --------------
.../Precursor/PrecursorEquilibrium.cu | 273 ------------------
.../Precursor/Precursor_Device.cuh | 51 ----
5 files changed, 613 deletions(-)
delete mode 100644 src/gpu/core/BoundaryConditions/Precursor/PrecursorDistributionsSubgridDistance.cu
delete mode 100644 src/gpu/core/BoundaryConditions/Precursor/PrecursorEquilibrium.cu
diff --git a/src/gpu/core/BoundaryConditions/Precursor/Precursor.cu b/src/gpu/core/BoundaryConditions/Precursor/Precursor.cu
index 8fd6f638f..1c9a047a1 100644
--- a/src/gpu/core/BoundaryConditions/Precursor/Precursor.cu
+++ b/src/gpu/core/BoundaryConditions/Precursor/Precursor.cu
@@ -77,44 +77,6 @@ void PrecursorNonReflectiveCompressible(
getLastCudaError("PrecursorNonReflectiveCompressible_Device execution failed");
}
-void PrecursorEquilibrium(
- LBMSimulationParameter* parameterDevice,
- QforPrecursorBoundaryConditions* boundaryCondition,
- real timeRatio,
- real velocityRatio)
-{
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
-
- PrecursorEquilibrium_Device<<< grid.grid, grid.threads >>>(
- boundaryCondition->k,
- boundaryCondition->numberOfBCnodes,
- boundaryCondition->numberOfPrecursorNodes,
- parameterDevice->omega,
- parameterDevice->distributions.f[0],
- parameterDevice->neighborX,
- parameterDevice->neighborX,
- parameterDevice->neighborX,
- boundaryCondition->planeNeighbor0PP,
- boundaryCondition->planeNeighbor0PM,
- boundaryCondition->planeNeighbor0MP,
- boundaryCondition->planeNeighbor0MM,
- boundaryCondition->weights0PP,
- boundaryCondition->weights0PM,
- boundaryCondition->weights0MP,
- boundaryCondition->weights0MM,
- boundaryCondition->last,
- boundaryCondition->current,
- boundaryCondition->velocityX,
- boundaryCondition->velocityY,
- boundaryCondition->velocityZ,
- timeRatio,
- velocityRatio,
- parameterDevice->numberOfNodes,
- parameterDevice->isEvenTimestep);
- getLastCudaError("PrecursorEquilibrium_Device execution failed");
-
-}
-
void PrecursorDistributions(
LBMSimulationParameter* parameterDevice,
QforPrecursorBoundaryConditions* boundaryCondition,
@@ -148,38 +110,3 @@ void PrecursorDistributions(
}
-void PrecursorDistributionsSubgridDistance(
- LBMSimulationParameter* parameterDevice,
- QforPrecursorBoundaryConditions* boundaryCondition,
- real timeRatio,
- real velocityRatio)
-{
-
- vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes);
-
- PrecursorDistributionsSubgridDistance_Device<<< grid.grid, grid.threads >>>(
- boundaryCondition->k,
- boundaryCondition->q27[0],
- boundaryCondition->sizeQ,
- boundaryCondition->numberOfBCnodes,
- boundaryCondition->numberOfPrecursorNodes,
- parameterDevice->distributions.f[0],
- parameterDevice->neighborX,
- parameterDevice->neighborY,
- parameterDevice->neighborZ,
- boundaryCondition->planeNeighbor0PP,
- boundaryCondition->planeNeighbor0PM,
- boundaryCondition->planeNeighbor0MP,
- boundaryCondition->planeNeighbor0MM,
- boundaryCondition->weights0PP,
- boundaryCondition->weights0PM,
- boundaryCondition->weights0MP,
- boundaryCondition->weights0MM,
- boundaryCondition->last,
- boundaryCondition->current,
- timeRatio,
- parameterDevice->numberOfNodes,
- parameterDevice->isEvenTimestep);
- getLastCudaError("PrecursorDistributionsSubgridDistance_Device execution failed");
-
-}
diff --git a/src/gpu/core/BoundaryConditions/Precursor/Precursor.h b/src/gpu/core/BoundaryConditions/Precursor/Precursor.h
index dcfc0f7ae..837da6953 100644
--- a/src/gpu/core/BoundaryConditions/Precursor/Precursor.h
+++ b/src/gpu/core/BoundaryConditions/Precursor/Precursor.h
@@ -38,10 +38,6 @@ class Parameter;
void PrecursorNonReflectiveCompressible(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio);
-void PrecursorEquilibrium(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio);
-
void PrecursorDistributions(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio);
-void PrecursorDistributionsSubgridDistance(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio);
-
#endif
diff --git a/src/gpu/core/BoundaryConditions/Precursor/PrecursorDistributionsSubgridDistance.cu b/src/gpu/core/BoundaryConditions/Precursor/PrecursorDistributionsSubgridDistance.cu
deleted file mode 100644
index 9eb78acb9..000000000
--- a/src/gpu/core/BoundaryConditions/Precursor/PrecursorDistributionsSubgridDistance.cu
+++ /dev/null
@@ -1,212 +0,0 @@
-//=======================================================================================
-// ____ ____ __ ______ __________ __ __ __ __
-// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
-// \ \ | | | | | |_) | | | | | | | / \ | |
-// \ \ | | | | | _ / | | | | | | / /\ \ | |
-// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
-// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
-// \ \ | | ________________________________________________________________
-// \ \ | | | ______________________________________________________________|
-// \ \| | | | __ __ __ __ ______ _______
-// \ | | |_____ | | | | | | | | | _ \ / _____)
-// \ | | _____| | | | | | | | | | | \ \ \_______
-// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
-// \ _____| |__| |________| \_______/ |__| |______/ (_______/
-//
-// 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 Henry Korb, Henrik Asmuth, Martin Schoenherr
-//======================================================================================
-#include "LBM/LB.h"
-#include <basics/constants/NumericConstants.h>
-#include <lbm/constants/D3Q27.h>
-#include <lbm/MacroscopicQuantities.h>
-
-#include "LBM/GPUHelperFunctions/KernelUtilities.h"
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-using namespace vf::gpu;
-
-////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
-// NOTE: Has not been tested after bug fix!
-__global__ void PrecursorDistributionsSubgridDistance_Device(
- int* subgridDistanceIndices,
- real* subgridDistances,
- int sizeQ,
- int numberOfBCnodes,
- int numberOfPrecursorNodes,
- real* distributions,
- uint* neighborX,
- uint* neighborY,
- uint* neighborZ,
- uint* neighbors0PP,
- uint* neighbors0PM,
- uint* neighbors0MP,
- uint* neighbors0MM,
- real* weights0PP,
- real* weights0PM,
- real* weights0MP,
- real* weights0MM,
- real* fsLast,
- real* fsNext,
- real timeRatio,
- unsigned long long numberOfLBnodes,
- bool isEvenTimestep)
-{
- ////////////////////////////////////////////////////////////////////////////////
- //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
- //!
- const unsigned nodeIndex = getNodeIndex();
-
- if(nodeIndex>=numberOfBCnodes) return;
-
- uint kNeighbor0PP = neighbors0PP[nodeIndex];
- real d0PP = weights0PP[nodeIndex];
-
- real f0LastInterp, f1LastInterp, f2LastInterp, f3LastInterp, f4LastInterp, f5LastInterp, f6LastInterp, f7LastInterp, f8LastInterp;
- real f0NextInterp, f1NextInterp, f2NextInterp, f3NextInterp, f4NextInterp, f5NextInterp, f6NextInterp, f7NextInterp, f8NextInterp;
-
- real* f0Last = fsLast;
- real* f1Last = &fsLast[ numberOfPrecursorNodes];
- real* f2Last = &fsLast[2*numberOfPrecursorNodes];
- real* f3Last = &fsLast[3*numberOfPrecursorNodes];
- real* f4Last = &fsLast[4*numberOfPrecursorNodes];
- real* f5Last = &fsLast[5*numberOfPrecursorNodes];
- real* f6Last = &fsLast[6*numberOfPrecursorNodes];
- real* f7Last = &fsLast[7*numberOfPrecursorNodes];
- real* f8Last = &fsLast[8*numberOfPrecursorNodes];
-
- real* f0Next = fsNext;
- real* f1Next = &fsNext[ numberOfPrecursorNodes];
- real* f2Next = &fsNext[2*numberOfPrecursorNodes];
- real* f3Next = &fsNext[3*numberOfPrecursorNodes];
- real* f4Next = &fsNext[4*numberOfPrecursorNodes];
- real* f5Next = &fsNext[5*numberOfPrecursorNodes];
- real* f6Next = &fsNext[6*numberOfPrecursorNodes];
- real* f7Next = &fsNext[7*numberOfPrecursorNodes];
- real* f8Next = &fsNext[8*numberOfPrecursorNodes];
-
-
- if(d0PP<1e6)
- {
- uint kNeighbor0PM = neighbors0PM[nodeIndex];
- uint kNeighbor0MP = neighbors0MP[nodeIndex];
- uint kNeighbor0MM = neighbors0MM[nodeIndex];
-
- real d0PM = weights0PM[nodeIndex];
- real d0MP = weights0MP[nodeIndex];
- real d0MM = weights0MM[nodeIndex];
-
- real invWeightSum = 1.f/(d0PP+d0PM+d0MP+d0MM);
-
- f0LastInterp = (f0Last[kNeighbor0PP]*d0PP + f0Last[kNeighbor0PM]*d0PM + f0Last[kNeighbor0MP]*d0MP + f0Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f0NextInterp = (f0Next[kNeighbor0PP]*d0PP + f0Next[kNeighbor0PM]*d0PM + f0Next[kNeighbor0MP]*d0MP + f0Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f1LastInterp = (f1Last[kNeighbor0PP]*d0PP + f1Last[kNeighbor0PM]*d0PM + f1Last[kNeighbor0MP]*d0MP + f1Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f1NextInterp = (f1Next[kNeighbor0PP]*d0PP + f1Next[kNeighbor0PM]*d0PM + f1Next[kNeighbor0MP]*d0MP + f1Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f2LastInterp = (f2Last[kNeighbor0PP]*d0PP + f2Last[kNeighbor0PM]*d0PM + f2Last[kNeighbor0MP]*d0MP + f2Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f2NextInterp = (f2Next[kNeighbor0PP]*d0PP + f2Next[kNeighbor0PM]*d0PM + f2Next[kNeighbor0MP]*d0MP + f2Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f3LastInterp = (f3Last[kNeighbor0PP]*d0PP + f3Last[kNeighbor0PM]*d0PM + f3Last[kNeighbor0MP]*d0MP + f3Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f3NextInterp = (f3Next[kNeighbor0PP]*d0PP + f3Next[kNeighbor0PM]*d0PM + f3Next[kNeighbor0MP]*d0MP + f3Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f4LastInterp = (f4Last[kNeighbor0PP]*d0PP + f4Last[kNeighbor0PM]*d0PM + f4Last[kNeighbor0MP]*d0MP + f4Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f4NextInterp = (f4Next[kNeighbor0PP]*d0PP + f4Next[kNeighbor0PM]*d0PM + f4Next[kNeighbor0MP]*d0MP + f4Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f5LastInterp = (f5Last[kNeighbor0PP]*d0PP + f5Last[kNeighbor0PM]*d0PM + f5Last[kNeighbor0MP]*d0MP + f5Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f5NextInterp = (f5Next[kNeighbor0PP]*d0PP + f5Next[kNeighbor0PM]*d0PM + f5Next[kNeighbor0MP]*d0MP + f5Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f6LastInterp = (f6Last[kNeighbor0PP]*d0PP + f6Last[kNeighbor0PM]*d0PM + f6Last[kNeighbor0MP]*d0MP + f6Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f6NextInterp = (f6Next[kNeighbor0PP]*d0PP + f6Next[kNeighbor0PM]*d0PM + f6Next[kNeighbor0MP]*d0MP + f6Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f7LastInterp = (f7Last[kNeighbor0PP]*d0PP + f7Last[kNeighbor0PM]*d0PM + f7Last[kNeighbor0MP]*d0MP + f7Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f7NextInterp = (f7Next[kNeighbor0PP]*d0PP + f7Next[kNeighbor0PM]*d0PM + f7Next[kNeighbor0MP]*d0MP + f7Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- f8LastInterp = (f8Last[kNeighbor0PP]*d0PP + f8Last[kNeighbor0PM]*d0PM + f8Last[kNeighbor0MP]*d0MP + f8Last[kNeighbor0MM]*d0MM)*invWeightSum;
- f8NextInterp = (f8Next[kNeighbor0PP]*d0PP + f8Next[kNeighbor0PM]*d0PM + f8Next[kNeighbor0MP]*d0MP + f8Next[kNeighbor0MM]*d0MM)*invWeightSum;
-
- } else {
- f0LastInterp = f0Last[kNeighbor0PP];
- f1LastInterp = f1Last[kNeighbor0PP];
- f2LastInterp = f2Last[kNeighbor0PP];
- f3LastInterp = f3Last[kNeighbor0PP];
- f4LastInterp = f4Last[kNeighbor0PP];
- f5LastInterp = f5Last[kNeighbor0PP];
- f6LastInterp = f6Last[kNeighbor0PP];
- f7LastInterp = f7Last[kNeighbor0PP];
- f8LastInterp = f8Last[kNeighbor0PP];
-
- f0NextInterp = f0Next[kNeighbor0PP];
- f1NextInterp = f1Next[kNeighbor0PP];
- f2NextInterp = f2Next[kNeighbor0PP];
- f3NextInterp = f3Next[kNeighbor0PP];
- f4NextInterp = f4Next[kNeighbor0PP];
- f5NextInterp = f5Next[kNeighbor0PP];
- f6NextInterp = f6Next[kNeighbor0PP];
- f7NextInterp = f7Next[kNeighbor0PP];
- f8NextInterp = f8Next[kNeighbor0PP];
- }
- //////////////////////////////////////////////////////////////////////////
- //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep
- //! is based on the esoteric twist algorithm \ref <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier
- //! et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
- //!
- Distributions27 dist;
- getPointersToDistributions(dist, distributions, numberOfLBnodes, !isEvenTimestep);
-
- unsigned int KQK = subgridDistanceIndices[nodeIndex];
- // unsigned int k000= KQK;
- unsigned int kP00 = KQK;
- // unsigned int kM00 = neighborX[KQK];
- // unsigned int k0P0 = KQK;
- unsigned int k0M0 = neighborY[KQK];
- // unsigned int k00P = KQK;
- unsigned int k00M = neighborZ[KQK];
- // unsigned int kMM0 = neighborY[kM00];
- unsigned int kPP0 = KQK;
- unsigned int kPM0 = k0M0;
- // unsigned int kMP0 = kM00;
- // unsigned int kM0M = neighborZ[kM00];
- unsigned int kP0P = KQK;
- unsigned int kP0M = k00M;
- // unsigned int kM0P = kM00;
- unsigned int k0MM = neighborZ[k0M0];
- // unsigned int k0PM = k00M;
- // unsigned int k0MP = k0M0;
- unsigned int kPMP = k0M0;
- // unsigned int kMPM = kM0M;
- // unsigned int kMPP = kM00;
- unsigned int kPMM = k0MM;
- // unsigned int kMMP = kMM0;
- unsigned int kPPM = k00M;
- unsigned int kPPP = KQK;
- // unsigned int kMMM = neighborZ[kMM0];
- SubgridDistances27 qs;
- getPointersToSubgridDistances(qs, subgridDistances, sizeQ);
-
- real q;
- q = qs.q[dP00][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dP00][kP00] = f0LastInterp*(1.f-timeRatio) + f0NextInterp*timeRatio;
- q = qs.q[dPP0][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPP0][kPP0] = f1LastInterp*(1.f-timeRatio) + f1NextInterp*timeRatio;
- q = qs.q[dPM0][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPM0][kPM0] = f2LastInterp*(1.f-timeRatio) + f2NextInterp*timeRatio;
- q = qs.q[dP0P][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dP0P][kP0P] = f3LastInterp*(1.f-timeRatio) + f3NextInterp*timeRatio;
- q = qs.q[dP0M][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dP0M][kP0M] = f4LastInterp*(1.f-timeRatio) + f4NextInterp*timeRatio;
- q = qs.q[dPPP][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPPP][kPPP] = f5LastInterp*(1.f-timeRatio) + f5NextInterp*timeRatio;
- q = qs.q[dPMP][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPMP][kPMP] = f6LastInterp*(1.f-timeRatio) + f6NextInterp*timeRatio;
- q = qs.q[dPPM][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPPM][kPPM] = f7LastInterp*(1.f-timeRatio) + f7NextInterp*timeRatio;
- q = qs.q[dPMM][nodeIndex]; if(q>= c0o1 && q <= c1o1) dist.f[dPMM][kPMM] = f8LastInterp*(1.f-timeRatio) + f8NextInterp*timeRatio;
-
-}
-////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/core/BoundaryConditions/Precursor/PrecursorEquilibrium.cu b/src/gpu/core/BoundaryConditions/Precursor/PrecursorEquilibrium.cu
deleted file mode 100644
index 2c82a4981..000000000
--- a/src/gpu/core/BoundaryConditions/Precursor/PrecursorEquilibrium.cu
+++ /dev/null
@@ -1,273 +0,0 @@
-//=======================================================================================
-// ____ ____ __ ______ __________ __ __ __ __
-// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
-// \ \ | | | | | |_) | | | | | | | / \ | |
-// \ \ | | | | | _ / | | | | | | / /\ \ | |
-// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
-// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
-// \ \ | | ________________________________________________________________
-// \ \ | | | ______________________________________________________________|
-// \ \| | | | __ __ __ __ ______ _______
-// \ | | |_____ | | | | | | | | | _ \ / _____)
-// \ | | _____| | | | | | | | | | | \ \ \_______
-// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
-// \ _____| |__| |________| \_______/ |__| |______/ (_______/
-//
-// 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 Henry Korb, Henrik Asmuth, Martin Schoenherr
-//======================================================================================
-#include "LBM/LB.h"
-#include <basics/constants/NumericConstants.h>
-#include <lbm/constants/D3Q27.h>
-#include <lbm/MacroscopicQuantities.h>
-
-#include "LBM/GPUHelperFunctions/KernelUtilities.h"
-
-using namespace vf::basics::constant;
-using namespace vf::lbm::dir;
-using namespace vf::gpu;
-
-__global__ void PrecursorEquilibrium_Device(
- int *subgridDistanceIndices,
- int numberOfBCnodes,
- int numberOfPrecursorNodes,
- real omega,
- real* distributions,
- uint* neighborX,
- uint* neighborY,
- uint* neighborZ,
- uint* neighbors0PP,
- uint* neighbors0PM,
- uint* neighbors0MP,
- uint* neighbors0MM,
- real* weights0PP,
- real* weights0PM,
- real* weights0MP,
- real* weights0MM,
- real* vLast,
- real* vCurrent,
- real velocityX,
- real velocityY,
- real velocityZ,
- real timeRatio,
- real velocityRatio,
- unsigned long long numberOfLBnodes,
- bool isEvenTimestep)
-{
- ////////////////////////////////////////////////////////////////////////////////
- //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
- //!
- const unsigned nodeIndex = getNodeIndex();
-
- if(nodeIndex>=numberOfBCnodes) return;
-
- ////////////////////////////////////////////////////////////////////////////////
- // interpolation of velocity
- real vxLastInterpd, vyLastInterpd, vzLastInterpd;
- real vxNextInterpd, vyNextInterpd, vzNextInterpd;
-
- uint kNeighbor0PP = neighbors0PP[nodeIndex];
- real d0PP = weights0PP[nodeIndex];
-
- real* vxLast = vLast;
- real* vyLast = &vLast[numberOfPrecursorNodes];
- real* vzLast = &vLast[2*numberOfPrecursorNodes];
-
- real* vxCurrent = vCurrent;
- real* vyCurrent = &vCurrent[numberOfPrecursorNodes];
- real* vzCurrent = &vCurrent[2*numberOfPrecursorNodes];
-
- if(d0PP < 1e6)
- {
- uint kNeighbor0PM = neighbors0PM[nodeIndex];
- uint kNeighbor0MP = neighbors0MP[nodeIndex];
- uint kNeighbor0MM = neighbors0MM[nodeIndex];
-
- real d0PM = weights0PM[nodeIndex];
- real d0MP = weights0MP[nodeIndex];
- real d0MM = weights0MM[nodeIndex];
-
- real invWeightSum = 1.f/(d0PP+d0PM+d0MP+d0MM);
-
- vxLastInterpd = (vxLast[kNeighbor0PP]*d0PP + vxLast[kNeighbor0PM]*d0PM + vxLast[kNeighbor0MP]*d0MP + vxLast[kNeighbor0MM]*d0MM)*invWeightSum;
- vyLastInterpd = (vyLast[kNeighbor0PP]*d0PP + vyLast[kNeighbor0PM]*d0PM + vyLast[kNeighbor0MP]*d0MP + vyLast[kNeighbor0MM]*d0MM)*invWeightSum;
- vzLastInterpd = (vzLast[kNeighbor0PP]*d0PP + vzLast[kNeighbor0PM]*d0PM + vzLast[kNeighbor0MP]*d0MP + vzLast[kNeighbor0MM]*d0MM)*invWeightSum;
-
- vxNextInterpd = (vxCurrent[kNeighbor0PP]*d0PP + vxCurrent[kNeighbor0PM]*d0PM + vxCurrent[kNeighbor0MP]*d0MP + vxCurrent[kNeighbor0MM]*d0MM)*invWeightSum;
- vyNextInterpd = (vyCurrent[kNeighbor0PP]*d0PP + vyCurrent[kNeighbor0PM]*d0PM + vyCurrent[kNeighbor0MP]*d0MP + vyCurrent[kNeighbor0MM]*d0MM)*invWeightSum;
- vzNextInterpd = (vzCurrent[kNeighbor0PP]*d0PP + vzCurrent[kNeighbor0PM]*d0PM + vzCurrent[kNeighbor0MP]*d0MP + vzCurrent[kNeighbor0MM]*d0MM)*invWeightSum;
- }
- else
- {
- vxLastInterpd = vxLast[kNeighbor0PP];
- vyLastInterpd = vyLast[kNeighbor0PP];
- vzLastInterpd = vzLast[kNeighbor0PP];
-
- vxNextInterpd = vxCurrent[kNeighbor0PP];
- vyNextInterpd = vyCurrent[kNeighbor0PP];
- vzNextInterpd = vzCurrent[kNeighbor0PP];
- }
-
- // if(k==16300) printf("%f %f %f\n", vxLastInterpd, vyLastInterpd, vzLastInterpd);
- real VeloX = (velocityX + (1.f-timeRatio)*vxLastInterpd + timeRatio*vxNextInterpd)/velocityRatio;
- real VeloY = (velocityY + (1.f-timeRatio)*vyLastInterpd + timeRatio*vyNextInterpd)/velocityRatio;
- real VeloZ = (velocityZ + (1.f-timeRatio)*vzLastInterpd + timeRatio*vzNextInterpd)/velocityRatio;
- // From here on just a copy of QVelDeviceCompZeroPress
- ////////////////////////////////////////////////////////////////////////////////
-
- //////////////////////////////////////////////////////////////////////////
- //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep
- //! is based on the esoteric twist algorithm \ref <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier
- //! et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
- //!
- Distributions27 dist;
- getPointersToDistributions(dist, distributions, numberOfLBnodes, !isEvenTimestep);
-
- unsigned int KQK = subgridDistanceIndices[nodeIndex]; //QK
- unsigned int k000 = KQK; //000
- unsigned int kP00 = KQK; //P00
- unsigned int kM00 = neighborX[KQK]; //M00
- unsigned int k0P0 = KQK; //n
- unsigned int k0M0 = neighborY[KQK]; //s
- unsigned int k00P = KQK; //t
- unsigned int k00M = neighborZ[KQK]; //b
- unsigned int kMM0 = neighborY[kM00]; //sw
- unsigned int kPP0 = KQK; //ne
- unsigned int kPM0 = k0M0; //se
- unsigned int kMP0 = kM00; //nw
- unsigned int kM0M = neighborZ[kM00]; //bw
- unsigned int kP0P = KQK; //te
- unsigned int kP0M = k00M; //be
- unsigned int k0PP = KQK; //tn
- unsigned int k0MM = neighborZ[k0M0]; //bs
- unsigned int kM0P = kM00; //tw
- unsigned int k0PM = k00M; //bn
- unsigned int k0MP = k0M0; //ts
- unsigned int kPMP = k0M0; //tse
- unsigned int kMPM = kM0M; //bnw
- unsigned int kMPP = kM00; //tnw
- unsigned int kPMM = k0MM; //bse
- unsigned int kMMP = kMM0; //tsw
- unsigned int kPPM = k00M; //bne
- unsigned int kPPP = KQK; //tne
- unsigned int kMMM = neighborZ[kMM0]; //bsw
-
- //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- // based on BGK Plus Comp
- //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- real f_M00 = (dist.f[dP00])[kP00];
- real f_P00 = (dist.f[dM00])[kM00];
- real f_0M0 = (dist.f[d0P0])[k0P0];
- real f_0P0 = (dist.f[d0M0])[k0M0];
- real f_00M = (dist.f[d00P])[k00P];
- real f_00P = (dist.f[d00M])[k00M];
- real f_MM0 = (dist.f[dPP0])[kPP0];
- real f_PP0 = (dist.f[dMM0])[kMM0];
- real f_MP0 = (dist.f[dPM0])[kPM0];
- real f_PM0 = (dist.f[dMP0])[kMP0];
- real f_M0M = (dist.f[dP0P])[kP0P];
- real f_P0P = (dist.f[dM0M])[kM0M];
- real f_M0P = (dist.f[dP0M])[kP0M];
- real f_P0M = (dist.f[dM0P])[kM0P];
- real f_0MM = (dist.f[vf::lbm::dir::d0PP])[k0PP];
- real f_0PP = (dist.f[d0MM])[k0MM];
- real f_0PM = (dist.f[d0MP])[k0MP];
- real f_0MP = (dist.f[d0PM])[k0PM];
- real f_000 = (dist.f[d000])[k000];
- real f_MMM = (dist.f[dPPP])[kPPP];
- real f_PPM = (dist.f[dMMP])[kMMP];
- real f_MPM = (dist.f[dPMP])[kPMP];
- real f_PMM = (dist.f[dMPP])[kMPP];
- real f_MMP = (dist.f[dPPM])[kPPM];
- real f_PPP = (dist.f[dMMM])[kMMM];
- real f_MPP = (dist.f[dPMM])[kPMM];
- real f_PMP = (dist.f[dMPM])[kMPM];
-
- ////////////////////////////////////////////////////////////////////////////////
- //! - Set macroscopic quantities
- //!
- real drho = c0o1;
-
- real vx1 = VeloX;
-
- real vx2 = VeloY;
-
- real vx3 = VeloZ;
-
- real cusq = c3o2 * (vx1 * vx1 + vx2 * vx2 + vx3 * vx3);
-
- ////////////////////////////////////////////////////////////////////////////////
- f_000 = c8o27* (drho-(drho+c1o1)*cusq);
- f_P00 = c2o27* (drho+(drho+c1o1)*(c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cusq));
- f_M00 = c2o27* (drho+(drho+c1o1)*(c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cusq));
- f_0P0 = c2o27* (drho+(drho+c1o1)*(c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cusq));
- f_0M0 = c2o27* (drho+(drho+c1o1)*(c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cusq));
- f_00P = c2o27* (drho+(drho+c1o1)*(c3o1*( vx3)+c9o2*( vx3)*( vx3)-cusq));
- f_00M = c2o27* (drho+(drho+c1o1)*(c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cusq));
- f_PP0 = c1o54* (drho+(drho+c1o1)*(c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cusq));
- f_MM0 = c1o54* (drho+(drho+c1o1)*(c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cusq));
- f_PM0 = c1o54* (drho+(drho+c1o1)*(c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cusq));
- f_MP0 = c1o54* (drho+(drho+c1o1)*(c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cusq));
- f_P0P = c1o54* (drho+(drho+c1o1)*(c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cusq));
- f_M0M = c1o54* (drho+(drho+c1o1)*(c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cusq));
- f_P0M = c1o54* (drho+(drho+c1o1)*(c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cusq));
- f_M0P = c1o54* (drho+(drho+c1o1)*(c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cusq));
- f_0PP = c1o54* (drho+(drho+c1o1)*(c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cusq));
- f_0MM = c1o54* (drho+(drho+c1o1)*(c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cusq));
- f_0PM = c1o54* (drho+(drho+c1o1)*(c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cusq));
- f_0MP = c1o54* (drho+(drho+c1o1)*(c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cusq));
- f_PPP = c1o216*(drho+(drho+c1o1)*(c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cusq));
- f_MMM = c1o216*(drho+(drho+c1o1)*(c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cusq));
- f_PPM = c1o216*(drho+(drho+c1o1)*(c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cusq));
- f_MMP = c1o216*(drho+(drho+c1o1)*(c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cusq));
- f_PMP = c1o216*(drho+(drho+c1o1)*(c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cusq));
- f_MPM = c1o216*(drho+(drho+c1o1)*(c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cusq));
- f_PMM = c1o216*(drho+(drho+c1o1)*(c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cusq));
- f_MPP = c1o216*(drho+(drho+c1o1)*(c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cusq));
-
- ////////////////////////////////////////////////////////////////////////////////
- //! write the new distributions to the bc nodes
- //!
- (dist.f[dP00])[kP00] = f_M00;
- (dist.f[dPP0])[kPP0] = f_MM0;
- (dist.f[dP0M])[kP0M] = f_M0P;
- (dist.f[dPM0])[kPM0] = f_MP0;
- (dist.f[dPMP])[kPMP] = f_MPM;
- (dist.f[dP0P])[kP0P] = f_M0M;
- (dist.f[dPPM])[kPPM] = f_MMP;
- (dist.f[dPPP])[kPPP] = f_MMM;
- (dist.f[dPMM])[kPMM] = f_MPP;
-
- (dist.f[dM00])[kM00] = f_P00;
- (dist.f[dMM0])[kMM0] = f_PP0;
- (dist.f[dM0M])[kM0M] = f_P0P;
- (dist.f[dMP0])[kMP0] = f_PM0;
- (dist.f[dM0P])[kM0P] = f_P0M;
- (dist.f[dMMM])[kMMM] = f_PPP;
- (dist.f[dMMP])[kMMP] = f_PPM;
- (dist.f[dMPP])[kMPP] = f_PMM;
- (dist.f[dMPM])[kMPM] = f_PMP;
-
- (dist.f[d0P0])[k0P0] = f_0M0;
- (dist.f[d0M0])[k0M0] = f_0P0;
- (dist.f[d00P])[k00P] = f_00M;
- (dist.f[d00M])[k00M] = f_00P;
- (dist.f[vf::lbm::dir::d0PP])[k0PP] = f_0MM;
- (dist.f[d0MM])[k0MM] = f_0PP;
- (dist.f[d0PM])[k0PM] = f_0MP;
- (dist.f[d0MP])[k0MP] = f_0PM;
- (dist.f[d000])[k000] = f_000;
-}
-
diff --git a/src/gpu/core/BoundaryConditions/Precursor/Precursor_Device.cuh b/src/gpu/core/BoundaryConditions/Precursor/Precursor_Device.cuh
index e22b78811..df7ab901e 100644
--- a/src/gpu/core/BoundaryConditions/Precursor/Precursor_Device.cuh
+++ b/src/gpu/core/BoundaryConditions/Precursor/Precursor_Device.cuh
@@ -62,33 +62,6 @@ __global__ void PrecursorNonReflectiveCompressible_Device(
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
-__global__ void PrecursorEquilibrium_Device(
- int* subgridDistanceIndices,
- int numberOfBCnodes,
- int numberOfPrecursorNodes,
- real omega,
- real* distributions,
- uint* neighborX,
- uint* neighborY,
- uint* neighborZ,
- uint* neighborsNT,
- uint* neighborsNB,
- uint* neighborsST,
- uint* neighborsSB,
- real* weights0PP,
- real* weights0PM,
- real* weights0MP,
- real* weights0MM,
- real* vLast,
- real* vCurrent,
- real velocityX,
- real velocityY,
- real velocityZ,
- real timeRatio,
- real velocityRatio,
- unsigned long long numberOfLBnodes,
- bool isEvenTimestep);
-
__global__ void PrecursorDistributions_Device(
int* subgridDistanceIndices,
int numberOfBCNodes,
@@ -111,28 +84,4 @@ __global__ void PrecursorDistributions_Device(
unsigned long long numberOfLBnodes,
bool isEvenTimestep);
-__global__ void PrecursorDistributionsSubgridDistance_Device(
- int* subgridDistanceIndices,
- real* subgridDistances,
- int sizeQ,
- int numberOfBCNodes,
- int numberOfPrecursorNodes,
- real* distributions,
- uint* neighborX,
- uint* neighborY,
- uint* neighborZ,
- uint* neighborsNT,
- uint* neighborsNB,
- uint* neighborsST,
- uint* neighborsSB,
- real* weights0PP,
- real* weights0PM,
- real* weights0MP,
- real* weights0MM,
- real* fsLast,
- real* fsNext,
- real timeRatio,
- unsigned long long numberOfLBnodes,
- bool isEvenTimestep);
-
#endif
--
GitLab