VirtualFluids 0.2.0
Parallel CFD LBM Solver
Loading...
Searching...
No Matches
NoSlipBounceBack.cu
Go to the documentation of this file.
1//=======================================================================================
2// ____ ____ __ ______ __________ __ __ __ __
3// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
4// \ \ | | | | | |_) | | | | | | | / \ | |
5// \ \ | | | | | _ / | | | | | | / /\ \ | |
6// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
7// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
8// \ \ | | ________________________________________________________________
9// \ \ | | | ______________________________________________________________|
10// \ \| | | | __ __ __ __ ______ _______
11// \ | | |_____ | | | | | | | | | _ \ / _____)
12// \ | | _____| | | | | | | | | | | \ \ \_______
13// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
14// \ _____| |__| |________| \_______/ |__| |______/ (_______/
15//
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.
20//
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
24// for more details.
25//
26// SPDX-License-Identifier: GPL-3.0-or-later
27// SPDX-FileCopyrightText: Copyright © VirtualFluids Project contributors, see AUTHORS.md in root folder
28//
29//! \addtogroup gpu_BoundaryConditions BoundaryConditions
30//! \ingroup gpu_core core
31//! \{
32//! \author Martin Schoenherr, Anna Wellmann
33//======================================================================================
34#include "Calculation/Calculation.h"
35#include "lbm/constants/D3Q27.h"
36#include <basics/constants/NumericConstants.h>
37#include "Utilities/KernelUtilities.h"
38#include "cuda_helper/CudaIndexCalculation.h"
39
40using namespace vf::basics::constant;
41using namespace vf::lbm::dir;
42namespace vf::gpu {
43
44__global__ void NoSlipBounceBack_Device(
45 real* distributions,
46 int* subgridDistanceIndices,
47 real* subgridDistances,
48 unsigned int numberOfBCnodes,
49 unsigned int* neighborX,
50 unsigned int* neighborY,
51 unsigned int* neighborZ,
52 unsigned long long numberOfLBnodes,
53 bool isEvenTimestep)
54{
55 //////////////////////////////////////////////////////////////////////////
56 //! The no-slip boundary condition is executed in the following steps
57 //!
58 ////////////////////////////////////////////////////////////////////////////////
59 //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
60 //!
61 const unsigned nodeIndex = vf::cuda::get1DIndexFrom2DBlock();
62
63 //////////////////////////////////////////////////////////////////////////
64 // run for all indices in size of boundary condition (numberOfBCnodes)
65 if(nodeIndex < numberOfBCnodes)
66 {
67 //////////////////////////////////////////////////////////////////////////
68 //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on timestep is based on the esoteric twist algorithm \ref
69 //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a>
70 //!
71 Distributions27 dist;
72 getPointersToDistributions(dist, distributions, numberOfLBnodes, isEvenTimestep);
73
74 ////////////////////////////////////////////////////////////////////////////////
75 //! - Set local subgrid distances (q's)
76 //!
77 SubgridDistances27 subgridD;
78 getPointersToSubgridDistances(subgridD, subgridDistances, numberOfBCnodes);
79
80 ////////////////////////////////////////////////////////////////////////////////
81 //! - Set neighbor indices (necessary for indirect addressing)
82 //!
83 unsigned int indexOfBCnode = subgridDistanceIndices[nodeIndex];
84 unsigned int ke = indexOfBCnode;
85 unsigned int kw = neighborX[indexOfBCnode];
86 unsigned int kn = indexOfBCnode;
87 unsigned int ks = neighborY[indexOfBCnode];
88 unsigned int kt = indexOfBCnode;
89 unsigned int kb = neighborZ[indexOfBCnode];
90 unsigned int ksw = neighborY[kw];
91 unsigned int kne = indexOfBCnode;
92 unsigned int kse = ks;
93 unsigned int knw = kw;
94 unsigned int kbw = neighborZ[kw];
95 unsigned int kte = indexOfBCnode;
96 unsigned int kbe = kb;
97 unsigned int ktw = kw;
98 unsigned int kbs = neighborZ[ks];
99 unsigned int ktn = indexOfBCnode;
100 unsigned int kbn = kb;
101 unsigned int kts = ks;
102 unsigned int ktse = ks;
103 unsigned int kbnw = kbw;
104 unsigned int ktnw = kw;
105 unsigned int kbse = kbs;
106 unsigned int ktsw = ksw;
107 unsigned int kbne = kb;
108 unsigned int ktne = indexOfBCnode;
109 unsigned int kbsw = neighborZ[ksw];
110
111 ////////////////////////////////////////////////////////////////////////////////
112 //! - Set local distributions
113 //!
114 real f_W = (dist.f[dP00])[ke ];
115 real f_E = (dist.f[dM00])[kw ];
116 real f_S = (dist.f[d0P0])[kn ];
117 real f_N = (dist.f[d0M0])[ks ];
118 real f_B = (dist.f[d00P])[kt ];
119 real f_T = (dist.f[d00M])[kb ];
120 real f_SW = (dist.f[dPP0])[kne ];
121 real f_NE = (dist.f[dMM0])[ksw ];
122 real f_NW = (dist.f[dPM0])[kse ];
123 real f_SE = (dist.f[dMP0])[knw ];
124 real f_BW = (dist.f[dP0P])[kte ];
125 real f_TE = (dist.f[dM0M])[kbw ];
126 real f_TW = (dist.f[dP0M])[kbe ];
127 real f_BE = (dist.f[dM0P])[ktw ];
128 real f_BS = (dist.f[d0PP])[ktn ];
129 real f_TN = (dist.f[d0MM])[kbs ];
130 real f_TS = (dist.f[d0PM])[kbn ];
131 real f_BN = (dist.f[d0MP])[kts ];
132 real f_BSW = (dist.f[dPPP])[ktne ];
133 real f_BNE = (dist.f[dMMP])[ktsw ];
134 real f_BNW = (dist.f[dPMP])[ktse ];
135 real f_BSE = (dist.f[dMPP])[ktnw ];
136 real f_TSW = (dist.f[dPPM])[kbne ];
137 real f_TNE = (dist.f[dMMM])[kbsw ];
138 real f_TNW = (dist.f[dPMM])[kbse ];
139 real f_TSE = (dist.f[dMPM])[kbnw ];
140
141 ////////////////////////////////////////////////////////////////////////////////
142 //! - change the pointer to write the results in the correct array
143 //!
144 getPointersToDistributions(dist, distributions, numberOfLBnodes, !isEvenTimestep);
145
146 ////////////////////////////////////////////////////////////////////////////////
147 //! - rewrite distributions if there is a sub-grid distance (q) in same direction
148 real q;
149 q = (subgridD.q[dP00])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dM00])[kw ]=f_E ;
150 q = (subgridD.q[dM00])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dP00])[ke ]=f_W ;
151 q = (subgridD.q[d0P0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0M0])[ks ]=f_N ;
152 q = (subgridD.q[d0M0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0P0])[kn ]=f_S ;
153 q = (subgridD.q[d00P])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d00M])[kb ]=f_T ;
154 q = (subgridD.q[d00M])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d00P])[kt ]=f_B ;
155 q = (subgridD.q[dPP0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMM0])[ksw ]=f_NE ;
156 q = (subgridD.q[dMM0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPP0])[kne ]=f_SW ;
157 q = (subgridD.q[dPM0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMP0])[knw ]=f_SE ;
158 q = (subgridD.q[dMP0])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPM0])[kse ]=f_NW ;
159 q = (subgridD.q[dP0P])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dM0M])[kbw ]=f_TE ;
160 q = (subgridD.q[dM0M])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dP0P])[kte ]=f_BW ;
161 q = (subgridD.q[dP0M])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dM0P])[ktw ]=f_BE ;
162 q = (subgridD.q[dM0P])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dP0M])[kbe ]=f_TW ;
163 q = (subgridD.q[d0PP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0MM])[kbs ]=f_TN ;
164 q = (subgridD.q[d0MM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0PP])[ktn ]=f_BS ;
165 q = (subgridD.q[d0PM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0MP])[kts ]=f_BN ;
166 q = (subgridD.q[d0MP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[d0PM])[kbn ]=f_TS ;
167 q = (subgridD.q[dPPP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMMM])[kbsw]=f_TNE;
168 q = (subgridD.q[dMMM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPPP])[ktne]=f_BSW;
169 q = (subgridD.q[dPPM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMMP])[ktsw]=f_BNE;
170 q = (subgridD.q[dMMP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPPM])[kbne]=f_TSW;
171 q = (subgridD.q[dPMP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMPM])[kbnw]=f_TSE;
172 q = (subgridD.q[dMPM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPMP])[ktse]=f_BNW;
173 q = (subgridD.q[dPMM])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dMPP])[ktnw]=f_BSE;
174 q = (subgridD.q[dMPP])[nodeIndex]; if (q>=c0o1 && q<=c1o1) (dist.f[dPMM])[kbse]=f_TNW;
175 }
176}
177
178}
179
180//! \}