From b5349b8b80c752d0d74e750a74b828d4fb72f0da Mon Sep 17 00:00:00 2001 From: Anna Wellmann <a.wellmann@tu-braunschweig.de> Date: Thu, 7 Jul 2022 14:28:20 +0200 Subject: [PATCH] Rename D3Q27 directions (cherry picked from commit 587f612efe06940eed8c8e15cf7457668e380ef8) --- .../GridReaderGenerator/GridGenerator.cpp | 54 +- .../GPU/AdvectionDiffusion27chim.cu | 364 ++++++------ .../GPU/AdvectionDiffusionBCs27.cu | 532 +++++++++--------- src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu | 156 ++--- src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu | 324 +++++------ .../VirtualFluids_GPU/GPU/Cumulant27chim.cu | 324 +++++------ src/gpu/VirtualFluids_GPU/GPU/Init27.cu | 162 +++--- .../GPU/InitAdvectionDiffusion27.cu | 162 +++--- src/gpu/VirtualFluids_GPU/LBM/D3Q27.h | 56 +- 9 files changed, 1067 insertions(+), 1067 deletions(-) diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp index 7a89a5aa8..96d3ebcf7 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp @@ -152,33 +152,33 @@ void GridGenerator::allocArrays_BoundaryQs() real *QQ = para->getParH()->slipBC.q27[0]; unsigned int sizeQ = para->getParH()->numberOfSlipBCnodes; QforBoundaryConditions Q; - Q.q27[dirE] = &QQ[dirE * sizeQ]; - Q.q27[dirW] = &QQ[dirW * sizeQ]; - Q.q27[dirN] = &QQ[dirN * sizeQ]; - Q.q27[dirS] = &QQ[dirS * sizeQ]; - Q.q27[dirT] = &QQ[dirT * sizeQ]; - Q.q27[dirB] = &QQ[dirB * sizeQ]; - Q.q27[dirNE] = &QQ[dirNE * sizeQ]; - Q.q27[dirSW] = &QQ[dirSW * sizeQ]; - Q.q27[dirSE] = &QQ[dirSE * sizeQ]; - Q.q27[dirNW] = &QQ[dirNW * sizeQ]; - Q.q27[dirTE] = &QQ[dirTE * sizeQ]; - Q.q27[dirBW] = &QQ[dirBW * sizeQ]; - Q.q27[dirBE] = &QQ[dirBE * sizeQ]; - Q.q27[dirTW] = &QQ[dirTW * sizeQ]; - Q.q27[dirTN] = &QQ[dirTN * sizeQ]; - Q.q27[dirBS] = &QQ[dirBS * sizeQ]; - Q.q27[dirBN] = &QQ[dirBN * sizeQ]; - Q.q27[dirTS] = &QQ[dirTS * sizeQ]; - Q.q27[dirREST] = &QQ[dirREST * sizeQ]; - Q.q27[dirTNE] = &QQ[dirTNE * sizeQ]; - Q.q27[dirTSW] = &QQ[dirTSW * sizeQ]; - Q.q27[dirTSE] = &QQ[dirTSE * sizeQ]; - Q.q27[dirTNW] = &QQ[dirTNW * sizeQ]; - Q.q27[dirBNE] = &QQ[dirBNE * sizeQ]; - Q.q27[dirBSW] = &QQ[dirBSW * sizeQ]; - Q.q27[dirBSE] = &QQ[dirBSE * sizeQ]; - Q.q27[dirBNW] = &QQ[dirBNW * sizeQ]; + Q.q27[E] = &QQ[E * sizeQ]; + Q.q27[W] = &QQ[W * sizeQ]; + Q.q27[N] = &QQ[N * sizeQ]; + Q.q27[S] = &QQ[S * sizeQ]; + Q.q27[T] = &QQ[T * sizeQ]; + Q.q27[B] = &QQ[B * sizeQ]; + Q.q27[NE] = &QQ[NE * sizeQ]; + Q.q27[SW] = &QQ[SW * sizeQ]; + Q.q27[SE] = &QQ[SE * sizeQ]; + Q.q27[NW] = &QQ[NW * sizeQ]; + Q.q27[TE] = &QQ[TE * sizeQ]; + Q.q27[BW] = &QQ[BW * sizeQ]; + Q.q27[BE] = &QQ[BE * sizeQ]; + Q.q27[TW] = &QQ[TW * sizeQ]; + Q.q27[TN] = &QQ[TN * sizeQ]; + Q.q27[BS] = &QQ[BS * sizeQ]; + Q.q27[BN] = &QQ[BN * sizeQ]; + Q.q27[TS] = &QQ[TS * sizeQ]; + Q.q27[REST] = &QQ[REST * sizeQ]; + Q.q27[TNE] = &QQ[TNE * sizeQ]; + Q.q27[TSW] = &QQ[TSW * sizeQ]; + Q.q27[TSE] = &QQ[TSE * sizeQ]; + Q.q27[TNW] = &QQ[TNW * sizeQ]; + Q.q27[BNE] = &QQ[BNE * sizeQ]; + Q.q27[BSW] = &QQ[BSW * sizeQ]; + Q.q27[BSE] = &QQ[BSE * sizeQ]; + Q.q27[BNW] = &QQ[BNW * sizeQ]; builder->getSlipQs(Q.q27, 0); diff --git a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu index 8b99bd07f..39e9a5916 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu @@ -109,125 +109,125 @@ extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device Distributions27 dist; if (isEvenTimestep) { - dist.f[dirE ] = &distributions[dirE *size_Mat]; - dist.f[dirW ] = &distributions[dirW *size_Mat]; - dist.f[dirN ] = &distributions[dirN *size_Mat]; - dist.f[dirS ] = &distributions[dirS *size_Mat]; - dist.f[dirT ] = &distributions[dirT *size_Mat]; - dist.f[dirB ] = &distributions[dirB *size_Mat]; - dist.f[dirNE ] = &distributions[dirNE *size_Mat]; - dist.f[dirSW ] = &distributions[dirSW *size_Mat]; - dist.f[dirSE ] = &distributions[dirSE *size_Mat]; - dist.f[dirNW ] = &distributions[dirNW *size_Mat]; - dist.f[dirTE ] = &distributions[dirTE *size_Mat]; - dist.f[dirBW ] = &distributions[dirBW *size_Mat]; - dist.f[dirBE ] = &distributions[dirBE *size_Mat]; - dist.f[dirTW ] = &distributions[dirTW *size_Mat]; - dist.f[dirTN ] = &distributions[dirTN *size_Mat]; - dist.f[dirBS ] = &distributions[dirBS *size_Mat]; - dist.f[dirBN ] = &distributions[dirBN *size_Mat]; - dist.f[dirTS ] = &distributions[dirTS *size_Mat]; + dist.f[E ] = &distributions[E *size_Mat]; + dist.f[W ] = &distributions[W *size_Mat]; + dist.f[N ] = &distributions[N *size_Mat]; + dist.f[S ] = &distributions[S *size_Mat]; + dist.f[T ] = &distributions[T *size_Mat]; + dist.f[B ] = &distributions[B *size_Mat]; + dist.f[NE ] = &distributions[NE *size_Mat]; + dist.f[SW ] = &distributions[SW *size_Mat]; + dist.f[SE ] = &distributions[SE *size_Mat]; + dist.f[NW ] = &distributions[NW *size_Mat]; + dist.f[TE ] = &distributions[TE *size_Mat]; + dist.f[BW ] = &distributions[BW *size_Mat]; + dist.f[BE ] = &distributions[BE *size_Mat]; + dist.f[TW ] = &distributions[TW *size_Mat]; + dist.f[TN ] = &distributions[TN *size_Mat]; + dist.f[BS ] = &distributions[BS *size_Mat]; + dist.f[BN ] = &distributions[BN *size_Mat]; + dist.f[TS ] = &distributions[TS *size_Mat]; dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirTNE *size_Mat]; - dist.f[dirTSW ] = &distributions[dirTSW *size_Mat]; - dist.f[dirTSE ] = &distributions[dirTSE *size_Mat]; - dist.f[dirTNW ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNE ] = &distributions[dirBNE *size_Mat]; - dist.f[dirBSW ] = &distributions[dirBSW *size_Mat]; - dist.f[dirBSE ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNW ] = &distributions[dirBNW *size_Mat]; + dist.f[TNE ] = &distributions[TNE *size_Mat]; + dist.f[TSW ] = &distributions[TSW *size_Mat]; + dist.f[TSE ] = &distributions[TSE *size_Mat]; + dist.f[TNW ] = &distributions[TNW *size_Mat]; + dist.f[BNE ] = &distributions[BNE *size_Mat]; + dist.f[BSW ] = &distributions[BSW *size_Mat]; + dist.f[BSE ] = &distributions[BSE *size_Mat]; + dist.f[BNW ] = &distributions[BNW *size_Mat]; } else { - dist.f[dirW ] = &distributions[dirE *size_Mat]; - dist.f[dirE ] = &distributions[dirW *size_Mat]; - dist.f[dirS ] = &distributions[dirN *size_Mat]; - dist.f[dirN ] = &distributions[dirS *size_Mat]; - dist.f[dirB ] = &distributions[dirT *size_Mat]; - dist.f[dirT ] = &distributions[dirB *size_Mat]; - dist.f[dirSW ] = &distributions[dirNE *size_Mat]; - dist.f[dirNE ] = &distributions[dirSW *size_Mat]; - dist.f[dirNW ] = &distributions[dirSE *size_Mat]; - dist.f[dirSE ] = &distributions[dirNW *size_Mat]; - dist.f[dirBW ] = &distributions[dirTE *size_Mat]; - dist.f[dirTE ] = &distributions[dirBW *size_Mat]; - dist.f[dirTW ] = &distributions[dirBE *size_Mat]; - dist.f[dirBE ] = &distributions[dirTW *size_Mat]; - dist.f[dirBS ] = &distributions[dirTN *size_Mat]; - dist.f[dirTN ] = &distributions[dirBS *size_Mat]; - dist.f[dirTS ] = &distributions[dirBN *size_Mat]; - dist.f[dirBN ] = &distributions[dirTS *size_Mat]; + dist.f[W ] = &distributions[E *size_Mat]; + dist.f[E ] = &distributions[W *size_Mat]; + dist.f[S ] = &distributions[N *size_Mat]; + dist.f[N ] = &distributions[S *size_Mat]; + dist.f[B ] = &distributions[T *size_Mat]; + dist.f[T ] = &distributions[B *size_Mat]; + dist.f[SW ] = &distributions[NE *size_Mat]; + dist.f[NE ] = &distributions[SW *size_Mat]; + dist.f[NW ] = &distributions[SE *size_Mat]; + dist.f[SE ] = &distributions[NW *size_Mat]; + dist.f[BW ] = &distributions[TE *size_Mat]; + dist.f[TE ] = &distributions[BW *size_Mat]; + dist.f[TW ] = &distributions[BE *size_Mat]; + dist.f[BE ] = &distributions[TW *size_Mat]; + dist.f[BS ] = &distributions[TN *size_Mat]; + dist.f[TN ] = &distributions[BS *size_Mat]; + dist.f[TS ] = &distributions[BN *size_Mat]; + dist.f[BN ] = &distributions[TS *size_Mat]; dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirBSW ] = &distributions[dirTNE *size_Mat]; - dist.f[dirBNE ] = &distributions[dirTSW *size_Mat]; - dist.f[dirBNW ] = &distributions[dirTSE *size_Mat]; - dist.f[dirBSE ] = &distributions[dirTNW *size_Mat]; - dist.f[dirTSW ] = &distributions[dirBNE *size_Mat]; - dist.f[dirTNE ] = &distributions[dirBSW *size_Mat]; - dist.f[dirTNW ] = &distributions[dirBSE *size_Mat]; - dist.f[dirTSE ] = &distributions[dirBNW *size_Mat]; + dist.f[BSW ] = &distributions[TNE *size_Mat]; + dist.f[BNE ] = &distributions[TSW *size_Mat]; + dist.f[BNW ] = &distributions[TSE *size_Mat]; + dist.f[BSE ] = &distributions[TNW *size_Mat]; + dist.f[TSW ] = &distributions[BNE *size_Mat]; + dist.f[TNE ] = &distributions[BSW *size_Mat]; + dist.f[TNW ] = &distributions[BSE *size_Mat]; + dist.f[TSE ] = &distributions[BNW *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// Distributions27 distAD; if (isEvenTimestep) { - distAD.f[dirE ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirW ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirTS *size_Mat]; + distAD.f[E ] = &distributionsAD[E *size_Mat]; + distAD.f[W ] = &distributionsAD[W *size_Mat]; + distAD.f[N ] = &distributionsAD[N *size_Mat]; + distAD.f[S ] = &distributionsAD[S *size_Mat]; + distAD.f[T ] = &distributionsAD[T *size_Mat]; + distAD.f[B ] = &distributionsAD[B *size_Mat]; + distAD.f[NE ] = &distributionsAD[NE *size_Mat]; + distAD.f[SW ] = &distributionsAD[SW *size_Mat]; + distAD.f[SE ] = &distributionsAD[SE *size_Mat]; + distAD.f[NW ] = &distributionsAD[NW *size_Mat]; + distAD.f[TE ] = &distributionsAD[TE *size_Mat]; + distAD.f[BW ] = &distributionsAD[BW *size_Mat]; + distAD.f[BE ] = &distributionsAD[BE *size_Mat]; + distAD.f[TW ] = &distributionsAD[TW *size_Mat]; + distAD.f[TN ] = &distributionsAD[TN *size_Mat]; + distAD.f[BS ] = &distributionsAD[BS *size_Mat]; + distAD.f[BN ] = &distributionsAD[BN *size_Mat]; + distAD.f[TS ] = &distributionsAD[TS *size_Mat]; distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[TNE ] = &distributionsAD[TNE *size_Mat]; + distAD.f[TSW ] = &distributionsAD[TSW *size_Mat]; + distAD.f[TSE ] = &distributionsAD[TSE *size_Mat]; + distAD.f[TNW ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNE ] = &distributionsAD[BNE *size_Mat]; + distAD.f[BSW ] = &distributionsAD[BSW *size_Mat]; + distAD.f[BSE ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNW ] = &distributionsAD[BNW *size_Mat]; } else { - distAD.f[dirW ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirE ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirTS *size_Mat]; + distAD.f[W ] = &distributionsAD[E *size_Mat]; + distAD.f[E ] = &distributionsAD[W *size_Mat]; + distAD.f[S ] = &distributionsAD[N *size_Mat]; + distAD.f[N ] = &distributionsAD[S *size_Mat]; + distAD.f[B ] = &distributionsAD[T *size_Mat]; + distAD.f[T ] = &distributionsAD[B *size_Mat]; + distAD.f[SW ] = &distributionsAD[NE *size_Mat]; + distAD.f[NE ] = &distributionsAD[SW *size_Mat]; + distAD.f[NW ] = &distributionsAD[SE *size_Mat]; + distAD.f[SE ] = &distributionsAD[NW *size_Mat]; + distAD.f[BW ] = &distributionsAD[TE *size_Mat]; + distAD.f[TE ] = &distributionsAD[BW *size_Mat]; + distAD.f[TW ] = &distributionsAD[BE *size_Mat]; + distAD.f[BE ] = &distributionsAD[TW *size_Mat]; + distAD.f[BS ] = &distributionsAD[TN *size_Mat]; + distAD.f[TN ] = &distributionsAD[BS *size_Mat]; + distAD.f[TS ] = &distributionsAD[BN *size_Mat]; + distAD.f[BN ] = &distributionsAD[TS *size_Mat]; distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[BSW ] = &distributionsAD[TNE *size_Mat]; + distAD.f[BNE ] = &distributionsAD[TSW *size_Mat]; + distAD.f[BNW ] = &distributionsAD[TSE *size_Mat]; + distAD.f[BSE ] = &distributionsAD[TNW *size_Mat]; + distAD.f[TSW ] = &distributionsAD[BNE *size_Mat]; + distAD.f[TNE ] = &distributionsAD[BSW *size_Mat]; + distAD.f[TNW ] = &distributionsAD[BSE *size_Mat]; + distAD.f[TSE ] = &distributionsAD[BNW *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) @@ -241,63 +241,63 @@ extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device //////////////////////////////////////////////////////////////////////////////////// //! - Set local distributions Fluid //! - real fcbb = (dist.f[dirE ])[k]; - real fabb = (dist.f[dirW ])[kw]; - real fbcb = (dist.f[dirN ])[k]; - real fbab = (dist.f[dirS ])[ks]; - real fbbc = (dist.f[dirT ])[k]; - real fbba = (dist.f[dirB ])[kb]; - real fccb = (dist.f[dirNE ])[k]; - real faab = (dist.f[dirSW ])[ksw]; - real fcab = (dist.f[dirSE ])[ks]; - real facb = (dist.f[dirNW ])[kw]; - real fcbc = (dist.f[dirTE ])[k]; - real faba = (dist.f[dirBW ])[kbw]; - real fcba = (dist.f[dirBE ])[kb]; - real fabc = (dist.f[dirTW ])[kw]; - real fbcc = (dist.f[dirTN ])[k]; - real fbaa = (dist.f[dirBS ])[kbs]; - real fbca = (dist.f[dirBN ])[kb]; - real fbac = (dist.f[dirTS ])[ks]; + real fcbb = (dist.f[E ])[k]; + real fabb = (dist.f[W ])[kw]; + real fbcb = (dist.f[N ])[k]; + real fbab = (dist.f[S ])[ks]; + real fbbc = (dist.f[T ])[k]; + real fbba = (dist.f[B ])[kb]; + real fccb = (dist.f[NE ])[k]; + real faab = (dist.f[SW ])[ksw]; + real fcab = (dist.f[SE ])[ks]; + real facb = (dist.f[NW ])[kw]; + real fcbc = (dist.f[TE ])[k]; + real faba = (dist.f[BW ])[kbw]; + real fcba = (dist.f[BE ])[kb]; + real fabc = (dist.f[TW ])[kw]; + real fbcc = (dist.f[TN ])[k]; + real fbaa = (dist.f[BS ])[kbs]; + real fbca = (dist.f[BN ])[kb]; + real fbac = (dist.f[TS ])[ks]; real fbbb = (dist.f[dirREST])[k]; - real fccc = (dist.f[dirTNE ])[k]; - real faac = (dist.f[dirTSW ])[ksw]; - real fcac = (dist.f[dirTSE ])[ks]; - real facc = (dist.f[dirTNW ])[kw]; - real fcca = (dist.f[dirBNE ])[kb]; - real faaa = (dist.f[dirBSW ])[kbsw]; - real fcaa = (dist.f[dirBSE ])[kbs]; - real faca = (dist.f[dirBNW ])[kbw]; + real fccc = (dist.f[TNE ])[k]; + real faac = (dist.f[TSW ])[ksw]; + real fcac = (dist.f[TSE ])[ks]; + real facc = (dist.f[TNW ])[kw]; + real fcca = (dist.f[BNE ])[kb]; + real faaa = (dist.f[BSW ])[kbsw]; + real fcaa = (dist.f[BSE ])[kbs]; + real faca = (dist.f[BNW ])[kbw]; //////////////////////////////////////////////////////////////////////////////////// //! - Set local distributions Advection Diffusion //! - real mfcbb = (distAD.f[dirE ])[k]; - real mfabb = (distAD.f[dirW ])[kw]; - real mfbcb = (distAD.f[dirN ])[k]; - real mfbab = (distAD.f[dirS ])[ks]; - real mfbbc = (distAD.f[dirT ])[k]; - real mfbba = (distAD.f[dirB ])[kb]; - real mfccb = (distAD.f[dirNE ])[k]; - real mfaab = (distAD.f[dirSW ])[ksw]; - real mfcab = (distAD.f[dirSE ])[ks]; - real mfacb = (distAD.f[dirNW ])[kw]; - real mfcbc = (distAD.f[dirTE ])[k]; - real mfaba = (distAD.f[dirBW ])[kbw]; - real mfcba = (distAD.f[dirBE ])[kb]; - real mfabc = (distAD.f[dirTW ])[kw]; - real mfbcc = (distAD.f[dirTN ])[k]; - real mfbaa = (distAD.f[dirBS ])[kbs]; - real mfbca = (distAD.f[dirBN ])[kb]; - real mfbac = (distAD.f[dirTS ])[ks]; + real mfcbb = (distAD.f[E ])[k]; + real mfabb = (distAD.f[W ])[kw]; + real mfbcb = (distAD.f[N ])[k]; + real mfbab = (distAD.f[S ])[ks]; + real mfbbc = (distAD.f[T ])[k]; + real mfbba = (distAD.f[B ])[kb]; + real mfccb = (distAD.f[NE ])[k]; + real mfaab = (distAD.f[SW ])[ksw]; + real mfcab = (distAD.f[SE ])[ks]; + real mfacb = (distAD.f[NW ])[kw]; + real mfcbc = (distAD.f[TE ])[k]; + real mfaba = (distAD.f[BW ])[kbw]; + real mfcba = (distAD.f[BE ])[kb]; + real mfabc = (distAD.f[TW ])[kw]; + real mfbcc = (distAD.f[TN ])[k]; + real mfbaa = (distAD.f[BS ])[kbs]; + real mfbca = (distAD.f[BN ])[kb]; + real mfbac = (distAD.f[TS ])[ks]; real mfbbb = (distAD.f[dirREST])[k]; - real mfccc = (distAD.f[dirTNE ])[k]; - real mfaac = (distAD.f[dirTSW ])[ksw]; - real mfcac = (distAD.f[dirTSE ])[ks]; - real mfacc = (distAD.f[dirTNW ])[kw]; - real mfcca = (distAD.f[dirBNE ])[kb]; - real mfaaa = (distAD.f[dirBSW ])[kbsw]; - real mfcaa = (distAD.f[dirBSE ])[kbs]; - real mfaca = (distAD.f[dirBNW ])[kbw]; + real mfccc = (distAD.f[TNE ])[k]; + real mfaac = (distAD.f[TSW ])[ksw]; + real mfcac = (distAD.f[TSE ])[ks]; + real mfacc = (distAD.f[TNW ])[kw]; + real mfcca = (distAD.f[BNE ])[kb]; + real mfaaa = (distAD.f[BSW ])[kbsw]; + real mfcaa = (distAD.f[BSE ])[kbs]; + real mfaca = (distAD.f[BNW ])[kbw]; //////////////////////////////////////////////////////////////////////////////////// //! - Calculate density and velocity using pyramid summation for low round-off errors as in Eq. (J1)-(J3) \ref //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa.2015.05.001 ]</b></a> @@ -503,33 +503,33 @@ extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device //! stored arrays dependent on timestep is based on the esoteric twist algorithm //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a> //! - (distAD.f[dirE ])[k ] = mfabb; - (distAD.f[dirW ])[kw ] = mfcbb; - (distAD.f[dirN ])[k ] = mfbab; - (distAD.f[dirS ])[ks ] = mfbcb; - (distAD.f[dirT ])[k ] = mfbba; - (distAD.f[dirB ])[kb ] = mfbbc; - (distAD.f[dirNE ])[k ] = mfaab; - (distAD.f[dirSW ])[ksw ] = mfccb; - (distAD.f[dirSE ])[ks ] = mfacb; - (distAD.f[dirNW ])[kw ] = mfcab; - (distAD.f[dirTE ])[k ] = mfaba; - (distAD.f[dirBW ])[kbw ] = mfcbc; - (distAD.f[dirBE ])[kb ] = mfabc; - (distAD.f[dirTW ])[kw ] = mfcba; - (distAD.f[dirTN ])[k ] = mfbaa; - (distAD.f[dirBS ])[kbs ] = mfbcc; - (distAD.f[dirBN ])[kb ] = mfbac; - (distAD.f[dirTS ])[ks ] = mfbca; + (distAD.f[E ])[k ] = mfabb; + (distAD.f[W ])[kw ] = mfcbb; + (distAD.f[N ])[k ] = mfbab; + (distAD.f[S ])[ks ] = mfbcb; + (distAD.f[T ])[k ] = mfbba; + (distAD.f[B ])[kb ] = mfbbc; + (distAD.f[NE ])[k ] = mfaab; + (distAD.f[SW ])[ksw ] = mfccb; + (distAD.f[SE ])[ks ] = mfacb; + (distAD.f[NW ])[kw ] = mfcab; + (distAD.f[TE ])[k ] = mfaba; + (distAD.f[BW ])[kbw ] = mfcbc; + (distAD.f[BE ])[kb ] = mfabc; + (distAD.f[TW ])[kw ] = mfcba; + (distAD.f[TN ])[k ] = mfbaa; + (distAD.f[BS ])[kbs ] = mfbcc; + (distAD.f[BN ])[kb ] = mfbac; + (distAD.f[TS ])[ks ] = mfbca; (distAD.f[dirREST])[k ] = mfbbb; - (distAD.f[dirTNE ])[k ] = mfaaa; - (distAD.f[dirTSE ])[ks ] = mfaca; - (distAD.f[dirBNE ])[kb ] = mfaac; - (distAD.f[dirBSE ])[kbs ] = mfacc; - (distAD.f[dirTNW ])[kw ] = mfcaa; - (distAD.f[dirTSW ])[ksw ] = mfcca; - (distAD.f[dirBNW ])[kbw ] = mfcac; - (distAD.f[dirBSW ])[kbsw] = mfccc; + (distAD.f[TNE ])[k ] = mfaaa; + (distAD.f[TSE ])[ks ] = mfaca; + (distAD.f[BNE ])[kb ] = mfaac; + (distAD.f[BSE ])[kbs ] = mfacc; + (distAD.f[TNW ])[kw ] = mfcaa; + (distAD.f[TSW ])[ksw ] = mfcca; + (distAD.f[BNW ])[kbw ] = mfcac; + (distAD.f[BSW ])[kbsw] = mfccc; } } //////////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusionBCs27.cu b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusionBCs27.cu index abf8639d3..fff5b9f3f 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusionBCs27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusionBCs27.cu @@ -58,7 +58,7 @@ extern "C" __global__ void AD_SlipVelDeviceComp( real *distributionsAD, int *QindexArray, real *Qarrays, - uint numberOfQs, + uint numberBCnodes, real omegaDiffusivity, uint* neighborX, uint* neighborY, @@ -69,125 +69,125 @@ extern "C" __global__ void AD_SlipVelDeviceComp( Distributions27 D; if (isEvenTimestep) { - D.f[dirE ] = &distributions[dirE * size_Mat]; - D.f[dirW ] = &distributions[dirW * size_Mat]; - D.f[dirN ] = &distributions[dirN * size_Mat]; - D.f[dirS ] = &distributions[dirS * size_Mat]; - D.f[dirT ] = &distributions[dirT * size_Mat]; - D.f[dirB ] = &distributions[dirB * size_Mat]; - D.f[dirNE ] = &distributions[dirNE * size_Mat]; - D.f[dirSW ] = &distributions[dirSW * size_Mat]; - D.f[dirSE ] = &distributions[dirSE * size_Mat]; - D.f[dirNW ] = &distributions[dirNW * size_Mat]; - D.f[dirTE ] = &distributions[dirTE * size_Mat]; - D.f[dirBW ] = &distributions[dirBW * size_Mat]; - D.f[dirBE ] = &distributions[dirBE * size_Mat]; - D.f[dirTW ] = &distributions[dirTW * size_Mat]; - D.f[dirTN ] = &distributions[dirTN * size_Mat]; - D.f[dirBS ] = &distributions[dirBS * size_Mat]; - D.f[dirBN ] = &distributions[dirBN * size_Mat]; - D.f[dirTS ] = &distributions[dirTS * size_Mat]; - D.f[dirREST] = &distributions[dirREST * size_Mat]; - D.f[dirTNE ] = &distributions[dirTNE * size_Mat]; - D.f[dirTSW ] = &distributions[dirTSW * size_Mat]; - D.f[dirTSE ] = &distributions[dirTSE * size_Mat]; - D.f[dirTNW ] = &distributions[dirTNW * size_Mat]; - D.f[dirBNE ] = &distributions[dirBNE * size_Mat]; - D.f[dirBSW ] = &distributions[dirBSW * size_Mat]; - D.f[dirBSE ] = &distributions[dirBSE * size_Mat]; - D.f[dirBNW ] = &distributions[dirBNW * size_Mat]; + D.f[E ] = &distributions[E * size_Mat]; + D.f[W ] = &distributions[W * size_Mat]; + D.f[N ] = &distributions[N * size_Mat]; + D.f[S ] = &distributions[S * size_Mat]; + D.f[T ] = &distributions[T * size_Mat]; + D.f[B ] = &distributions[B * size_Mat]; + D.f[NE ] = &distributions[NE * size_Mat]; + D.f[SW ] = &distributions[SW * size_Mat]; + D.f[SE ] = &distributions[SE * size_Mat]; + D.f[NW ] = &distributions[NW * size_Mat]; + D.f[TE ] = &distributions[TE * size_Mat]; + D.f[BW ] = &distributions[BW * size_Mat]; + D.f[BE ] = &distributions[BE * size_Mat]; + D.f[TW ] = &distributions[TW * size_Mat]; + D.f[TN ] = &distributions[TN * size_Mat]; + D.f[BS ] = &distributions[BS * size_Mat]; + D.f[BN ] = &distributions[BN * size_Mat]; + D.f[TS ] = &distributions[TS * size_Mat]; + D.f[REST] = &distributions[REST * size_Mat]; + D.f[TNE ] = &distributions[TNE * size_Mat]; + D.f[TSW ] = &distributions[TSW * size_Mat]; + D.f[TSE ] = &distributions[TSE * size_Mat]; + D.f[TNW ] = &distributions[TNW * size_Mat]; + D.f[BNE ] = &distributions[BNE * size_Mat]; + D.f[BSW ] = &distributions[BSW * size_Mat]; + D.f[BSE ] = &distributions[BSE * size_Mat]; + D.f[BNW ] = &distributions[BNW * size_Mat]; } else { - D.f[dirW ] = &distributions[dirE * size_Mat]; - D.f[dirE ] = &distributions[dirW * size_Mat]; - D.f[dirS ] = &distributions[dirN * size_Mat]; - D.f[dirN ] = &distributions[dirS * size_Mat]; - D.f[dirB ] = &distributions[dirT * size_Mat]; - D.f[dirT ] = &distributions[dirB * size_Mat]; - D.f[dirSW ] = &distributions[dirNE * size_Mat]; - D.f[dirNE ] = &distributions[dirSW * size_Mat]; - D.f[dirNW ] = &distributions[dirSE * size_Mat]; - D.f[dirSE ] = &distributions[dirNW * size_Mat]; - D.f[dirBW ] = &distributions[dirTE * size_Mat]; - D.f[dirTE ] = &distributions[dirBW * size_Mat]; - D.f[dirTW ] = &distributions[dirBE * size_Mat]; - D.f[dirBE ] = &distributions[dirTW * size_Mat]; - D.f[dirBS ] = &distributions[dirTN * size_Mat]; - D.f[dirTN ] = &distributions[dirBS * size_Mat]; - D.f[dirTS ] = &distributions[dirBN * size_Mat]; - D.f[dirBN ] = &distributions[dirTS * size_Mat]; - D.f[dirREST] = &distributions[dirREST * size_Mat]; - D.f[dirTNE ] = &distributions[dirBSW * size_Mat]; - D.f[dirTSW ] = &distributions[dirBNE * size_Mat]; - D.f[dirTSE ] = &distributions[dirBNW * size_Mat]; - D.f[dirTNW ] = &distributions[dirBSE * size_Mat]; - D.f[dirBNE ] = &distributions[dirTSW * size_Mat]; - D.f[dirBSW ] = &distributions[dirTNE * size_Mat]; - D.f[dirBSE ] = &distributions[dirTNW * size_Mat]; - D.f[dirBNW ] = &distributions[dirTSE * size_Mat]; + D.f[W ] = &distributions[E * size_Mat]; + D.f[E ] = &distributions[W * size_Mat]; + D.f[S ] = &distributions[N * size_Mat]; + D.f[N ] = &distributions[S * size_Mat]; + D.f[B ] = &distributions[T * size_Mat]; + D.f[T ] = &distributions[B * size_Mat]; + D.f[SW ] = &distributions[NE * size_Mat]; + D.f[NE ] = &distributions[SW * size_Mat]; + D.f[NW ] = &distributions[SE * size_Mat]; + D.f[SE ] = &distributions[NW * size_Mat]; + D.f[BW ] = &distributions[TE * size_Mat]; + D.f[TE ] = &distributions[BW * size_Mat]; + D.f[TW ] = &distributions[BE * size_Mat]; + D.f[BE ] = &distributions[TW * size_Mat]; + D.f[BS ] = &distributions[TN * size_Mat]; + D.f[TN ] = &distributions[BS * size_Mat]; + D.f[TS ] = &distributions[BN * size_Mat]; + D.f[BN ] = &distributions[TS * size_Mat]; + D.f[REST] = &distributions[REST * size_Mat]; + D.f[TNE ] = &distributions[BSW * size_Mat]; + D.f[TSW ] = &distributions[BNE * size_Mat]; + D.f[TSE ] = &distributions[BNW * size_Mat]; + D.f[TNW ] = &distributions[BSE * size_Mat]; + D.f[BNE ] = &distributions[TSW * size_Mat]; + D.f[BSW ] = &distributions[TNE * size_Mat]; + D.f[BSE ] = &distributions[TNW * size_Mat]; + D.f[BNW ] = &distributions[TSE * size_Mat]; } //////////////////////////////////////////////////////////////////////////////// Distributions27 DAD; if (isEvenTimestep) { - DAD.f[dirE ] = &distributionsAD[dirE * size_Mat]; - DAD.f[dirW ] = &distributionsAD[dirW * size_Mat]; - DAD.f[dirN ] = &distributionsAD[dirN * size_Mat]; - DAD.f[dirS ] = &distributionsAD[dirS * size_Mat]; - DAD.f[dirT ] = &distributionsAD[dirT * size_Mat]; - DAD.f[dirB ] = &distributionsAD[dirB * size_Mat]; - DAD.f[dirNE ] = &distributionsAD[dirNE * size_Mat]; - DAD.f[dirSW ] = &distributionsAD[dirSW * size_Mat]; - DAD.f[dirSE ] = &distributionsAD[dirSE * size_Mat]; - DAD.f[dirNW ] = &distributionsAD[dirNW * size_Mat]; - DAD.f[dirTE ] = &distributionsAD[dirTE * size_Mat]; - DAD.f[dirBW ] = &distributionsAD[dirBW * size_Mat]; - DAD.f[dirBE ] = &distributionsAD[dirBE * size_Mat]; - DAD.f[dirTW ] = &distributionsAD[dirTW * size_Mat]; - DAD.f[dirTN ] = &distributionsAD[dirTN * size_Mat]; - DAD.f[dirBS ] = &distributionsAD[dirBS * size_Mat]; - DAD.f[dirBN ] = &distributionsAD[dirBN * size_Mat]; - DAD.f[dirTS ] = &distributionsAD[dirTS * size_Mat]; - DAD.f[dirREST] = &distributionsAD[dirREST * size_Mat]; - DAD.f[dirTNE ] = &distributionsAD[dirTNE * size_Mat]; - DAD.f[dirTSW ] = &distributionsAD[dirTSW * size_Mat]; - DAD.f[dirTSE ] = &distributionsAD[dirTSE * size_Mat]; - DAD.f[dirTNW ] = &distributionsAD[dirTNW * size_Mat]; - DAD.f[dirBNE ] = &distributionsAD[dirBNE * size_Mat]; - DAD.f[dirBSW ] = &distributionsAD[dirBSW * size_Mat]; - DAD.f[dirBSE ] = &distributionsAD[dirBSE * size_Mat]; - DAD.f[dirBNW ] = &distributionsAD[dirBNW * size_Mat]; + DAD.f[E ] = &distributionsAD[E * size_Mat]; + DAD.f[W ] = &distributionsAD[W * size_Mat]; + DAD.f[N ] = &distributionsAD[N * size_Mat]; + DAD.f[S ] = &distributionsAD[S * size_Mat]; + DAD.f[T ] = &distributionsAD[T * size_Mat]; + DAD.f[B ] = &distributionsAD[B * size_Mat]; + DAD.f[NE ] = &distributionsAD[NE * size_Mat]; + DAD.f[SW ] = &distributionsAD[SW * size_Mat]; + DAD.f[SE ] = &distributionsAD[SE * size_Mat]; + DAD.f[NW ] = &distributionsAD[NW * size_Mat]; + DAD.f[TE ] = &distributionsAD[TE * size_Mat]; + DAD.f[BW ] = &distributionsAD[BW * size_Mat]; + DAD.f[BE ] = &distributionsAD[BE * size_Mat]; + DAD.f[TW ] = &distributionsAD[TW * size_Mat]; + DAD.f[TN ] = &distributionsAD[TN * size_Mat]; + DAD.f[BS ] = &distributionsAD[BS * size_Mat]; + DAD.f[BN ] = &distributionsAD[BN * size_Mat]; + DAD.f[TS ] = &distributionsAD[TS * size_Mat]; + DAD.f[REST] = &distributionsAD[REST * size_Mat]; + DAD.f[TNE ] = &distributionsAD[TNE * size_Mat]; + DAD.f[TSW ] = &distributionsAD[TSW * size_Mat]; + DAD.f[TSE ] = &distributionsAD[TSE * size_Mat]; + DAD.f[TNW ] = &distributionsAD[TNW * size_Mat]; + DAD.f[BNE ] = &distributionsAD[BNE * size_Mat]; + DAD.f[BSW ] = &distributionsAD[BSW * size_Mat]; + DAD.f[BSE ] = &distributionsAD[BSE * size_Mat]; + DAD.f[BNW ] = &distributionsAD[BNW * size_Mat]; } else { - DAD.f[dirW ] = &distributionsAD[dirE * size_Mat]; - DAD.f[dirE ] = &distributionsAD[dirW * size_Mat]; - DAD.f[dirS ] = &distributionsAD[dirN * size_Mat]; - DAD.f[dirN ] = &distributionsAD[dirS * size_Mat]; - DAD.f[dirB ] = &distributionsAD[dirT * size_Mat]; - DAD.f[dirT ] = &distributionsAD[dirB * size_Mat]; - DAD.f[dirSW ] = &distributionsAD[dirNE * size_Mat]; - DAD.f[dirNE ] = &distributionsAD[dirSW * size_Mat]; - DAD.f[dirNW ] = &distributionsAD[dirSE * size_Mat]; - DAD.f[dirSE ] = &distributionsAD[dirNW * size_Mat]; - DAD.f[dirBW ] = &distributionsAD[dirTE * size_Mat]; - DAD.f[dirTE ] = &distributionsAD[dirBW * size_Mat]; - DAD.f[dirTW ] = &distributionsAD[dirBE * size_Mat]; - DAD.f[dirBE ] = &distributionsAD[dirTW * size_Mat]; - DAD.f[dirBS ] = &distributionsAD[dirTN * size_Mat]; - DAD.f[dirTN ] = &distributionsAD[dirBS * size_Mat]; - DAD.f[dirTS ] = &distributionsAD[dirBN * size_Mat]; - DAD.f[dirBN ] = &distributionsAD[dirTS * size_Mat]; - DAD.f[dirREST] = &distributionsAD[dirREST * size_Mat]; - DAD.f[dirTNE ] = &distributionsAD[dirBSW * size_Mat]; - DAD.f[dirTSW ] = &distributionsAD[dirBNE * size_Mat]; - DAD.f[dirTSE ] = &distributionsAD[dirBNW * size_Mat]; - DAD.f[dirTNW ] = &distributionsAD[dirBSE * size_Mat]; - DAD.f[dirBNE ] = &distributionsAD[dirTSW * size_Mat]; - DAD.f[dirBSW ] = &distributionsAD[dirTNE * size_Mat]; - DAD.f[dirBSE ] = &distributionsAD[dirTNW * size_Mat]; - DAD.f[dirBNW ] = &distributionsAD[dirTSE * size_Mat]; + DAD.f[W ] = &distributionsAD[E * size_Mat]; + DAD.f[E ] = &distributionsAD[W * size_Mat]; + DAD.f[S ] = &distributionsAD[N * size_Mat]; + DAD.f[N ] = &distributionsAD[S * size_Mat]; + DAD.f[B ] = &distributionsAD[T * size_Mat]; + DAD.f[T ] = &distributionsAD[B * size_Mat]; + DAD.f[SW ] = &distributionsAD[NE * size_Mat]; + DAD.f[NE ] = &distributionsAD[SW * size_Mat]; + DAD.f[NW ] = &distributionsAD[SE * size_Mat]; + DAD.f[SE ] = &distributionsAD[NW * size_Mat]; + DAD.f[BW ] = &distributionsAD[TE * size_Mat]; + DAD.f[TE ] = &distributionsAD[BW * size_Mat]; + DAD.f[TW ] = &distributionsAD[BE * size_Mat]; + DAD.f[BE ] = &distributionsAD[TW * size_Mat]; + DAD.f[BS ] = &distributionsAD[TN * size_Mat]; + DAD.f[TN ] = &distributionsAD[BS * size_Mat]; + DAD.f[TS ] = &distributionsAD[BN * size_Mat]; + DAD.f[BN ] = &distributionsAD[TS * size_Mat]; + DAD.f[REST] = &distributionsAD[REST * size_Mat]; + DAD.f[TNE ] = &distributionsAD[BSW * size_Mat]; + DAD.f[TSW ] = &distributionsAD[BNE * size_Mat]; + DAD.f[TSE ] = &distributionsAD[BNW * size_Mat]; + DAD.f[TNW ] = &distributionsAD[BSE * size_Mat]; + DAD.f[BNE ] = &distributionsAD[TSW * size_Mat]; + DAD.f[BSW ] = &distributionsAD[TNE * size_Mat]; + DAD.f[BSE ] = &distributionsAD[TNW * size_Mat]; + DAD.f[BNW ] = &distributionsAD[TSE * size_Mat]; } //////////////////////////////////////////////////////////////////////////////// const unsigned x = threadIdx.x; // Globaler x-Index @@ -200,7 +200,7 @@ extern "C" __global__ void AD_SlipVelDeviceComp( const unsigned k = nx * (ny * z + y) + x; ////////////////////////////////////////////////////////////////////////// - if (k < numberOfQs) + if (k < numberBCnodes) { //////////////////////////////////////////////////////////////////////////////// real NormX = normalX[k]; @@ -212,32 +212,32 @@ extern "C" __global__ void AD_SlipVelDeviceComp( * 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 = &Qarrays[dirE * numberOfQs]; - q_dirW = &Qarrays[dirW * numberOfQs]; - q_dirN = &Qarrays[dirN * numberOfQs]; - q_dirS = &Qarrays[dirS * numberOfQs]; - q_dirT = &Qarrays[dirT * numberOfQs]; - q_dirB = &Qarrays[dirB * numberOfQs]; - q_dirNE = &Qarrays[dirNE * numberOfQs]; - q_dirSW = &Qarrays[dirSW * numberOfQs]; - q_dirSE = &Qarrays[dirSE * numberOfQs]; - q_dirNW = &Qarrays[dirNW * numberOfQs]; - q_dirTE = &Qarrays[dirTE * numberOfQs]; - q_dirBW = &Qarrays[dirBW * numberOfQs]; - q_dirBE = &Qarrays[dirBE * numberOfQs]; - q_dirTW = &Qarrays[dirTW * numberOfQs]; - q_dirTN = &Qarrays[dirTN * numberOfQs]; - q_dirBS = &Qarrays[dirBS * numberOfQs]; - q_dirBN = &Qarrays[dirBN * numberOfQs]; - q_dirTS = &Qarrays[dirTS * numberOfQs]; - q_dirTNE = &Qarrays[dirTNE * numberOfQs]; - q_dirTSW = &Qarrays[dirTSW * numberOfQs]; - q_dirTSE = &Qarrays[dirTSE * numberOfQs]; - q_dirTNW = &Qarrays[dirTNW * numberOfQs]; - q_dirBNE = &Qarrays[dirBNE * numberOfQs]; - q_dirBSW = &Qarrays[dirBSW * numberOfQs]; - q_dirBSE = &Qarrays[dirBSE * numberOfQs]; - q_dirBNW = &Qarrays[dirBNW * numberOfQs]; + q_dirE = &Qarrays[E * numberOfBCnodes]; + q_dirW = &Qarrays[W * numberOfBCnodes]; + q_dirN = &Qarrays[N * numberOfBCnodes]; + q_dirS = &Qarrays[S * numberOfBCnodes]; + q_dirT = &Qarrays[T * numberOfBCnodes]; + q_dirB = &Qarrays[B * numberOfBCnodes]; + q_dirNE = &Qarrays[NE * numberOfBCnodes]; + q_dirSW = &Qarrays[SW * numberOfBCnodes]; + q_dirSE = &Qarrays[SE * numberOfBCnodes]; + q_dirNW = &Qarrays[NW * numberOfBCnodes]; + q_dirTE = &Qarrays[TE * numberOfBCnodes]; + q_dirBW = &Qarrays[BW * numberOfBCnodes]; + q_dirBE = &Qarrays[BE * numberOfBCnodes]; + q_dirTW = &Qarrays[TW * numberOfBCnodes]; + q_dirTN = &Qarrays[TN * numberOfBCnodes]; + q_dirBS = &Qarrays[BS * numberOfBCnodes]; + q_dirBN = &Qarrays[BN * numberOfBCnodes]; + q_dirTS = &Qarrays[TS * numberOfBCnodes]; + q_dirTNE = &Qarrays[TNE * numberOfBCnodes]; + q_dirTSW = &Qarrays[TSW * numberOfBCnodes]; + q_dirTSE = &Qarrays[TSE * numberOfBCnodes]; + q_dirTNW = &Qarrays[TNW * numberOfBCnodes]; + q_dirBNE = &Qarrays[BNE * numberOfBCnodes]; + q_dirBSW = &Qarrays[BSW * numberOfBCnodes]; + q_dirBSE = &Qarrays[BSE * numberOfBCnodes]; + q_dirBNW = &Qarrays[BNW * numberOfBCnodes]; //////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = QindexArray[k]; @@ -272,32 +272,32 @@ extern "C" __global__ void AD_SlipVelDeviceComp( real f_E, f_W, f_N, f_S, f_T, f_B, f_NE, f_SW, f_SE, f_NW, f_TE, f_BW, f_BE, f_TW, f_TN, f_BS, f_BN, f_TS, f_TNE, f_TSW, f_TSE, f_TNW, f_BNE, f_BSW, f_BSE, f_BNW; - f_W = (D.f[dirE])[ke]; - f_E = (D.f[dirW])[kw]; - f_S = (D.f[dirN])[kn]; - f_N = (D.f[dirS])[ks]; - f_B = (D.f[dirT])[kt]; - f_T = (D.f[dirB])[kb]; - f_SW = (D.f[dirNE])[kne]; - f_NE = (D.f[dirSW])[ksw]; - f_NW = (D.f[dirSE])[kse]; - f_SE = (D.f[dirNW])[knw]; - f_BW = (D.f[dirTE])[kte]; - f_TE = (D.f[dirBW])[kbw]; - f_TW = (D.f[dirBE])[kbe]; - f_BE = (D.f[dirTW])[ktw]; - f_BS = (D.f[dirTN])[ktn]; - f_TN = (D.f[dirBS])[kbs]; - f_TS = (D.f[dirBN])[kbn]; - f_BN = (D.f[dirTS])[kts]; - f_BSW = (D.f[dirTNE])[ktne]; - f_BNE = (D.f[dirTSW])[ktsw]; - f_BNW = (D.f[dirTSE])[ktse]; - f_BSE = (D.f[dirTNW])[ktnw]; - f_TSW = (D.f[dirBNE])[kbne]; - f_TNE = (D.f[dirBSW])[kbsw]; - f_TNW = (D.f[dirBSE])[kbse]; - f_TSE = (D.f[dirBNW])[kbnw]; + f_W = (D.f[E])[ke]; + f_E = (D.f[W])[kw]; + f_S = (D.f[N])[kn]; + f_N = (D.f[S])[ks]; + f_B = (D.f[T])[kt]; + f_T = (D.f[B])[kb]; + f_SW = (D.f[NE])[kne]; + f_NE = (D.f[SW])[ksw]; + f_NW = (D.f[SE])[kse]; + f_SE = (D.f[NW])[knw]; + f_BW = (D.f[TE])[kte]; + f_TE = (D.f[BW])[kbw]; + f_TW = (D.f[BE])[kbe]; + f_BE = (D.f[TW])[ktw]; + f_BS = (D.f[TN])[ktn]; + f_TN = (D.f[BS])[kbs]; + f_TS = (D.f[BN])[kbn]; + f_BN = (D.f[TS])[kts]; + f_BSW = (D.f[TNE])[ktne]; + f_BNE = (D.f[TSW])[ktsw]; + f_BNW = (D.f[TSE])[ktse]; + f_BSE = (D.f[TNW])[ktnw]; + f_TSW = (D.f[BNE])[kbne]; + f_TNE = (D.f[BSW])[kbsw]; + f_TNW = (D.f[BSE])[kbse]; + f_TSE = (D.f[BNW])[kbnw]; //////////////////////////////////////////////////////////////////////////////// real vx1, vx2, vx3, drho, q; drho = f_TSE + f_TNW + f_TNE + f_TSW + f_BSE + f_BNW + f_BNE + f_BSW + @@ -320,92 +320,92 @@ extern "C" __global__ void AD_SlipVelDeviceComp( real cu_sq = c3o2 * (vx1 * vx1 + vx2 * vx2 + vx3 * vx3) * (c1o1 + drho); //////////////////////////////////////////////////////////////////////////////// - f_W = (DAD.f[dirE])[ke]; - f_E = (DAD.f[dirW])[kw]; - f_S = (DAD.f[dirN])[kn]; - f_N = (DAD.f[dirS])[ks]; - f_B = (DAD.f[dirT])[kt]; - f_T = (DAD.f[dirB])[kb]; - f_SW = (DAD.f[dirNE])[kne]; - f_NE = (DAD.f[dirSW])[ksw]; - f_NW = (DAD.f[dirSE])[kse]; - f_SE = (DAD.f[dirNW])[knw]; - f_BW = (DAD.f[dirTE])[kte]; - f_TE = (DAD.f[dirBW])[kbw]; - f_TW = (DAD.f[dirBE])[kbe]; - f_BE = (DAD.f[dirTW])[ktw]; - f_BS = (DAD.f[dirTN])[ktn]; - f_TN = (DAD.f[dirBS])[kbs]; - f_TS = (DAD.f[dirBN])[kbn]; - f_BN = (DAD.f[dirTS])[kts]; - f_BSW = (DAD.f[dirTNE])[ktne]; - f_BNE = (DAD.f[dirTSW])[ktsw]; - f_BNW = (DAD.f[dirTSE])[ktse]; - f_BSE = (DAD.f[dirTNW])[ktnw]; - f_TSW = (DAD.f[dirBNE])[kbne]; - f_TNE = (DAD.f[dirBSW])[kbsw]; - f_TNW = (DAD.f[dirBSE])[kbse]; - f_TSE = (DAD.f[dirBNW])[kbnw]; + f_W = (DAD.f[E])[ke]; + f_E = (DAD.f[W])[kw]; + f_S = (DAD.f[N])[kn]; + f_N = (DAD.f[S])[ks]; + f_B = (DAD.f[T])[kt]; + f_T = (DAD.f[B])[kb]; + f_SW = (DAD.f[NE])[kne]; + f_NE = (DAD.f[SW])[ksw]; + f_NW = (DAD.f[SE])[kse]; + f_SE = (DAD.f[NW])[knw]; + f_BW = (DAD.f[TE])[kte]; + f_TE = (DAD.f[BW])[kbw]; + f_TW = (DAD.f[BE])[kbe]; + f_BE = (DAD.f[TW])[ktw]; + f_BS = (DAD.f[TN])[ktn]; + f_TN = (DAD.f[BS])[kbs]; + f_TS = (DAD.f[BN])[kbn]; + f_BN = (DAD.f[TS])[kts]; + f_BSW = (DAD.f[TNE])[ktne]; + f_BNE = (DAD.f[TSW])[ktsw]; + f_BNW = (DAD.f[TSE])[ktse]; + f_BSE = (DAD.f[TNW])[ktnw]; + f_TSW = (DAD.f[BNE])[kbne]; + f_TNE = (DAD.f[BSW])[kbsw]; + f_TNW = (DAD.f[BSE])[kbse]; + f_TSE = (DAD.f[BNW])[kbnw]; ////////////////////////////////////////////////////////////////////////// if (!isEvenTimestep) { - DAD.f[dirE ] = &distributionsAD[dirE * size_Mat]; - DAD.f[dirW ] = &distributionsAD[dirW * size_Mat]; - DAD.f[dirN ] = &distributionsAD[dirN * size_Mat]; - DAD.f[dirS ] = &distributionsAD[dirS * size_Mat]; - DAD.f[dirT ] = &distributionsAD[dirT * size_Mat]; - DAD.f[dirB ] = &distributionsAD[dirB * size_Mat]; - DAD.f[dirNE ] = &distributionsAD[dirNE * size_Mat]; - DAD.f[dirSW ] = &distributionsAD[dirSW * size_Mat]; - DAD.f[dirSE ] = &distributionsAD[dirSE * size_Mat]; - DAD.f[dirNW ] = &distributionsAD[dirNW * size_Mat]; - DAD.f[dirTE ] = &distributionsAD[dirTE * size_Mat]; - DAD.f[dirBW ] = &distributionsAD[dirBW * size_Mat]; - DAD.f[dirBE ] = &distributionsAD[dirBE * size_Mat]; - DAD.f[dirTW ] = &distributionsAD[dirTW * size_Mat]; - DAD.f[dirTN ] = &distributionsAD[dirTN * size_Mat]; - DAD.f[dirBS ] = &distributionsAD[dirBS * size_Mat]; - DAD.f[dirBN ] = &distributionsAD[dirBN * size_Mat]; - DAD.f[dirTS ] = &distributionsAD[dirTS * size_Mat]; + DAD.f[E ] = &distributionsAD[E * size_Mat]; + DAD.f[W ] = &distributionsAD[W * size_Mat]; + DAD.f[N ] = &distributionsAD[N * size_Mat]; + DAD.f[S ] = &distributionsAD[S * size_Mat]; + DAD.f[T ] = &distributionsAD[T * size_Mat]; + DAD.f[B ] = &distributionsAD[B * size_Mat]; + DAD.f[NE ] = &distributionsAD[NE * size_Mat]; + DAD.f[SW ] = &distributionsAD[SW * size_Mat]; + DAD.f[SE ] = &distributionsAD[SE * size_Mat]; + DAD.f[NW ] = &distributionsAD[NW * size_Mat]; + DAD.f[TE ] = &distributionsAD[TE * size_Mat]; + DAD.f[BW ] = &distributionsAD[BW * size_Mat]; + DAD.f[BE ] = &distributionsAD[BE * size_Mat]; + DAD.f[TW ] = &distributionsAD[TW * size_Mat]; + DAD.f[TN ] = &distributionsAD[TN * size_Mat]; + DAD.f[BS ] = &distributionsAD[BS * size_Mat]; + DAD.f[BN ] = &distributionsAD[BN * size_Mat]; + DAD.f[TS ] = &distributionsAD[TS * size_Mat]; DAD.f[dirREST] = &distributionsAD[dirREST * size_Mat]; - DAD.f[dirTNE ] = &distributionsAD[dirTNE * size_Mat]; - DAD.f[dirTSW ] = &distributionsAD[dirTSW * size_Mat]; - DAD.f[dirTSE ] = &distributionsAD[dirTSE * size_Mat]; - DAD.f[dirTNW ] = &distributionsAD[dirTNW * size_Mat]; - DAD.f[dirBNE ] = &distributionsAD[dirBNE * size_Mat]; - DAD.f[dirBSW ] = &distributionsAD[dirBSW * size_Mat]; - DAD.f[dirBSE ] = &distributionsAD[dirBSE * size_Mat]; - DAD.f[dirBNW ] = &distributionsAD[dirBNW * size_Mat]; + DAD.f[TNE ] = &distributionsAD[TNE * size_Mat]; + DAD.f[TSW ] = &distributionsAD[TSW * size_Mat]; + DAD.f[TSE ] = &distributionsAD[TSE * size_Mat]; + DAD.f[TNW ] = &distributionsAD[TNW * size_Mat]; + DAD.f[BNE ] = &distributionsAD[BNE * size_Mat]; + DAD.f[BSW ] = &distributionsAD[BSW * size_Mat]; + DAD.f[BSE ] = &distributionsAD[BSE * size_Mat]; + DAD.f[BNW ] = &distributionsAD[BNW * size_Mat]; } else { - DAD.f[dirW ] = &distributionsAD[dirE * size_Mat]; - DAD.f[dirE ] = &distributionsAD[dirW * size_Mat]; - DAD.f[dirS ] = &distributionsAD[dirN * size_Mat]; - DAD.f[dirN ] = &distributionsAD[dirS * size_Mat]; - DAD.f[dirB ] = &distributionsAD[dirT * size_Mat]; - DAD.f[dirT ] = &distributionsAD[dirB * size_Mat]; - DAD.f[dirSW ] = &distributionsAD[dirNE * size_Mat]; - DAD.f[dirNE ] = &distributionsAD[dirSW * size_Mat]; - DAD.f[dirNW ] = &distributionsAD[dirSE * size_Mat]; - DAD.f[dirSE ] = &distributionsAD[dirNW * size_Mat]; - DAD.f[dirBW ] = &distributionsAD[dirTE * size_Mat]; - DAD.f[dirTE ] = &distributionsAD[dirBW * size_Mat]; - DAD.f[dirTW ] = &distributionsAD[dirBE * size_Mat]; - DAD.f[dirBE ] = &distributionsAD[dirTW * size_Mat]; - DAD.f[dirBS ] = &distributionsAD[dirTN * size_Mat]; - DAD.f[dirTN ] = &distributionsAD[dirBS * size_Mat]; - DAD.f[dirTS ] = &distributionsAD[dirBN * size_Mat]; - DAD.f[dirBN ] = &distributionsAD[dirTS * size_Mat]; + DAD.f[W ] = &distributionsAD[E * size_Mat]; + DAD.f[E ] = &distributionsAD[W * size_Mat]; + DAD.f[S ] = &distributionsAD[N * size_Mat]; + DAD.f[N ] = &distributionsAD[S * size_Mat]; + DAD.f[B ] = &distributionsAD[T * size_Mat]; + DAD.f[T ] = &distributionsAD[B * size_Mat]; + DAD.f[SW ] = &distributionsAD[NE * size_Mat]; + DAD.f[NE ] = &distributionsAD[SW * size_Mat]; + DAD.f[NW ] = &distributionsAD[SE * size_Mat]; + DAD.f[SE ] = &distributionsAD[NW * size_Mat]; + DAD.f[BW ] = &distributionsAD[TE * size_Mat]; + DAD.f[TE ] = &distributionsAD[BW * size_Mat]; + DAD.f[TW ] = &distributionsAD[BE * size_Mat]; + DAD.f[BE ] = &distributionsAD[TW * size_Mat]; + DAD.f[BS ] = &distributionsAD[TN * size_Mat]; + DAD.f[TN ] = &distributionsAD[BS * size_Mat]; + DAD.f[TS ] = &distributionsAD[BN * size_Mat]; + DAD.f[BN ] = &distributionsAD[TS * size_Mat]; DAD.f[dirREST] = &distributionsAD[dirREST * size_Mat]; - DAD.f[dirTNE ] = &distributionsAD[dirBSW * size_Mat]; - DAD.f[dirTSW ] = &distributionsAD[dirBNE * size_Mat]; - DAD.f[dirTSE ] = &distributionsAD[dirBNW * size_Mat]; - DAD.f[dirTNW ] = &distributionsAD[dirBSE * size_Mat]; - DAD.f[dirBNE ] = &distributionsAD[dirTSW * size_Mat]; - DAD.f[dirBSW ] = &distributionsAD[dirTNE * size_Mat]; - DAD.f[dirBSE ] = &distributionsAD[dirTNW * size_Mat]; - DAD.f[dirBNW ] = &distributionsAD[dirTSE * size_Mat]; + DAD.f[TNE ] = &distributionsAD[BSW * size_Mat]; + DAD.f[TSW ] = &distributionsAD[BNE * size_Mat]; + DAD.f[TSE ] = &distributionsAD[BNW * size_Mat]; + DAD.f[TNW ] = &distributionsAD[BSE * size_Mat]; + DAD.f[BNE ] = &distributionsAD[TSW * size_Mat]; + DAD.f[BSW ] = &distributionsAD[TNE * size_Mat]; + DAD.f[BSE ] = &distributionsAD[TNW * size_Mat]; + DAD.f[BNW ] = &distributionsAD[TSE * size_Mat]; } //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// real concentration = @@ -439,32 +439,32 @@ extern "C" __global__ void AD_SlipVelDeviceComp( real jTan3 = jx3 - NormJ * NormZ; //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - q = q_dirE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirW ])[kw ] = calcDistributionBC_AD(q, c2o27, vx1, cu_sq, f_E, f_W, omegaDiffusivity, jTan1, concentration); } - q = q_dirW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirE ])[ke ] = calcDistributionBC_AD(q, c2o27, -vx1, cu_sq, f_W, f_E, omegaDiffusivity, -jTan1, concentration); } - q = q_dirN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirS ])[ks ] = calcDistributionBC_AD(q, c2o27, vx2, cu_sq, f_N, f_S, omegaDiffusivity, jTan2, concentration); } - q = q_dirS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirN ])[kn ] = calcDistributionBC_AD(q, c2o27, -vx2, cu_sq, f_S, f_N, omegaDiffusivity, -jTan2, concentration); } - q = q_dirT[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirB ])[kb ] = calcDistributionBC_AD(q, c2o27, vx3, cu_sq, f_T, f_B, omegaDiffusivity, jTan3, concentration); } - q = q_dirB[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirT ])[kt ] = calcDistributionBC_AD(q, c2o27, -vx3, cu_sq, f_B, f_T, omegaDiffusivity, -jTan3, concentration); } - q = q_dirNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirSW ])[ksw ] = calcDistributionBC_AD(q, c1o54, vx1+vx2, cu_sq, f_NE, f_SW, omegaDiffusivity, jTan1+jTan2, concentration); } - q = q_dirSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirNE ])[kne ] = calcDistributionBC_AD(q, c1o54, -vx1-vx2, cu_sq, f_SW, f_NE, omegaDiffusivity, -jTan1-jTan2, concentration); } - q = q_dirSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirNW ])[knw ] = calcDistributionBC_AD(q, c1o54, vx1-vx2, cu_sq, f_SE, f_NW, omegaDiffusivity, jTan1-jTan2, concentration); } - q = q_dirNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirSE ])[kse ] = calcDistributionBC_AD(q, c1o54, -vx1+vx2, cu_sq, f_NW, f_SE, omegaDiffusivity, -jTan1+jTan2, concentration); } - q = q_dirTE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBW ])[kbw ] = calcDistributionBC_AD(q, c1o54, vx1 +vx3, cu_sq, f_TE, f_BW, omegaDiffusivity, jTan1 +jTan3, concentration); } - q = q_dirBW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTE ])[kte ] = calcDistributionBC_AD(q, c1o54, -vx1 -vx3, cu_sq, f_BW, f_TE, omegaDiffusivity, -jTan1 -jTan3, concentration); } - q = q_dirBE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTW ])[ktw ] = calcDistributionBC_AD(q, c1o54, vx1 -vx3, cu_sq, f_BE, f_TW, omegaDiffusivity, jTan1 -jTan3, concentration); } - q = q_dirTW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBE ])[kbe ] = calcDistributionBC_AD(q, c1o54, -vx1 +vx3, cu_sq, f_TW, f_BE, omegaDiffusivity, -jTan1 +jTan3, concentration); } - q = q_dirTN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBS ])[kbs ] = calcDistributionBC_AD(q, c1o54, vx2+vx3, cu_sq, f_TN, f_BS, omegaDiffusivity, jTan2+jTan3, concentration); } - q = q_dirBS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTN ])[ktn ] = calcDistributionBC_AD(q, c1o54, -vx2-vx3, cu_sq, f_BS, f_TN, omegaDiffusivity, -jTan2-jTan3, concentration); } - q = q_dirBN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTS ])[kts ] = calcDistributionBC_AD(q, c1o54, vx2-vx3, cu_sq, f_BN, f_TS, omegaDiffusivity, jTan2-jTan3, concentration); } - q = q_dirTS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBN ])[kbn ] = calcDistributionBC_AD(q, c1o54, -vx2+vx3, cu_sq, f_TS, f_BN, omegaDiffusivity, -jTan2+jTan3, concentration); } - q = q_dirTNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBSW])[kbsw] = calcDistributionBC_AD(q, c1o216, vx1+vx2+vx3, cu_sq, f_TNE, f_BSW, omegaDiffusivity, jTan1+jTan2+jTan3, concentration); } - q = q_dirBSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTNE])[ktne] = calcDistributionBC_AD(q, c1o216, -vx1-vx2-vx3, cu_sq, f_BSW, f_TNE, omegaDiffusivity, -jTan1-jTan2-jTan3, concentration); } - q = q_dirBNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTSW])[ktsw] = calcDistributionBC_AD(q, c1o216, vx1+vx2-vx3, cu_sq, f_BNE, f_TSW, omegaDiffusivity, jTan1+jTan2-jTan3, concentration); } - q = q_dirTSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBNE])[kbne] = calcDistributionBC_AD(q, c1o216, -vx1-vx2+vx3, cu_sq, f_TSW, f_BNE, omegaDiffusivity, -jTan1-jTan2+jTan3, concentration); } - q = q_dirTSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBNW])[kbnw] = calcDistributionBC_AD(q, c1o216, vx1-vx2+vx3, cu_sq, f_TSE, f_BNW, omegaDiffusivity, jTan1-jTan2+jTan3, concentration); } - q = q_dirBNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTSE])[ktse] = calcDistributionBC_AD(q, c1o216, -vx1+vx2-vx3, cu_sq, f_BNW, f_TSE, omegaDiffusivity, -jTan1+jTan2-jTan3, concentration); } - q = q_dirBSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirTNW])[ktnw] = calcDistributionBC_AD(q, c1o216, vx1-vx2-vx3, cu_sq, f_BSE, f_TNW, omegaDiffusivity, jTan1-jTan2-jTan3, concentration); } - q = q_dirTNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[dirBSE])[kbse] = calcDistributionBC_AD(q, c1o216, -vx1+vx2+vx3, cu_sq, f_TNW, f_BSE, omegaDiffusivity, -jTan1+jTan2+jTan3, concentration); } + q = q_dirE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[W ])[kw ] = calcDistributionBC_AD(q, c2o27, vx1, cu_sq, f_E, f_W, omegaDiffusivity, jTan1, concentration); } + q = q_dirW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[E ])[ke ] = calcDistributionBC_AD(q, c2o27, -vx1, cu_sq, f_W, f_E, omegaDiffusivity, -jTan1, concentration); } + q = q_dirN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[S ])[ks ] = calcDistributionBC_AD(q, c2o27, vx2, cu_sq, f_N, f_S, omegaDiffusivity, jTan2, concentration); } + q = q_dirS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[N ])[kn ] = calcDistributionBC_AD(q, c2o27, -vx2, cu_sq, f_S, f_N, omegaDiffusivity, -jTan2, concentration); } + q = q_dirT[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[B ])[kb ] = calcDistributionBC_AD(q, c2o27, vx3, cu_sq, f_T, f_B, omegaDiffusivity, jTan3, concentration); } + q = q_dirB[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[T ])[kt ] = calcDistributionBC_AD(q, c2o27, -vx3, cu_sq, f_B, f_T, omegaDiffusivity, -jTan3, concentration); } + q = q_dirNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[SW ])[ksw ] = calcDistributionBC_AD(q, c1o54, vx1+vx2, cu_sq, f_NE, f_SW, omegaDiffusivity, jTan1+jTan2, concentration); } + q = q_dirSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[NE ])[kne ] = calcDistributionBC_AD(q, c1o54, -vx1-vx2, cu_sq, f_SW, f_NE, omegaDiffusivity, -jTan1-jTan2, concentration); } + q = q_dirSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[NW ])[knw ] = calcDistributionBC_AD(q, c1o54, vx1-vx2, cu_sq, f_SE, f_NW, omegaDiffusivity, jTan1-jTan2, concentration); } + q = q_dirNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[SE ])[kse ] = calcDistributionBC_AD(q, c1o54, -vx1+vx2, cu_sq, f_NW, f_SE, omegaDiffusivity, -jTan1+jTan2, concentration); } + q = q_dirTE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BW ])[kbw ] = calcDistributionBC_AD(q, c1o54, vx1 +vx3, cu_sq, f_TE, f_BW, omegaDiffusivity, jTan1 +jTan3, concentration); } + q = q_dirBW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TE ])[kte ] = calcDistributionBC_AD(q, c1o54, -vx1 -vx3, cu_sq, f_BW, f_TE, omegaDiffusivity, -jTan1 -jTan3, concentration); } + q = q_dirBE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TW ])[ktw ] = calcDistributionBC_AD(q, c1o54, vx1 -vx3, cu_sq, f_BE, f_TW, omegaDiffusivity, jTan1 -jTan3, concentration); } + q = q_dirTW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BE ])[kbe ] = calcDistributionBC_AD(q, c1o54, -vx1 +vx3, cu_sq, f_TW, f_BE, omegaDiffusivity, -jTan1 +jTan3, concentration); } + q = q_dirTN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BS ])[kbs ] = calcDistributionBC_AD(q, c1o54, vx2+vx3, cu_sq, f_TN, f_BS, omegaDiffusivity, jTan2+jTan3, concentration); } + q = q_dirBS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TN ])[ktn ] = calcDistributionBC_AD(q, c1o54, -vx2-vx3, cu_sq, f_BS, f_TN, omegaDiffusivity, -jTan2-jTan3, concentration); } + q = q_dirBN[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TS ])[kts ] = calcDistributionBC_AD(q, c1o54, vx2-vx3, cu_sq, f_BN, f_TS, omegaDiffusivity, jTan2-jTan3, concentration); } + q = q_dirTS[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BN ])[kbn ] = calcDistributionBC_AD(q, c1o54, -vx2+vx3, cu_sq, f_TS, f_BN, omegaDiffusivity, -jTan2+jTan3, concentration); } + q = q_dirTNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BSW])[kbsw] = calcDistributionBC_AD(q, c1o216, vx1+vx2+vx3, cu_sq, f_TNE, f_BSW, omegaDiffusivity, jTan1+jTan2+jTan3, concentration); } + q = q_dirBSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TNE])[ktne] = calcDistributionBC_AD(q, c1o216, -vx1-vx2-vx3, cu_sq, f_BSW, f_TNE, omegaDiffusivity, -jTan1-jTan2-jTan3, concentration); } + q = q_dirBNE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TSW])[ktsw] = calcDistributionBC_AD(q, c1o216, vx1+vx2-vx3, cu_sq, f_BNE, f_TSW, omegaDiffusivity, jTan1+jTan2-jTan3, concentration); } + q = q_dirTSW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BNE])[kbne] = calcDistributionBC_AD(q, c1o216, -vx1-vx2+vx3, cu_sq, f_TSW, f_BNE, omegaDiffusivity, -jTan1-jTan2+jTan3, concentration); } + q = q_dirTSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BNW])[kbnw] = calcDistributionBC_AD(q, c1o216, vx1-vx2+vx3, cu_sq, f_TSE, f_BNW, omegaDiffusivity, jTan1-jTan2+jTan3, concentration); } + q = q_dirBNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TSE])[ktse] = calcDistributionBC_AD(q, c1o216, -vx1+vx2-vx3, cu_sq, f_BNW, f_TSE, omegaDiffusivity, -jTan1+jTan2-jTan3, concentration); } + q = q_dirBSE[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[TNW])[ktnw] = calcDistributionBC_AD(q, c1o216, vx1-vx2-vx3, cu_sq, f_BSE, f_TNW, omegaDiffusivity, jTan1-jTan2-jTan3, concentration); } + q = q_dirTNW[k]; if (q >= c0o1 && q <= c1o1) { (DAD.f[BSE])[kbse] = calcDistributionBC_AD(q, c1o216, -vx1+vx2+vx3, cu_sq, f_TNW, f_BSE, omegaDiffusivity, -jTan1+jTan2+jTan3, concentration); } //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// } } diff --git a/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu b/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu index 328579769..556771c1a 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu @@ -75,63 +75,63 @@ extern "C" __global__ void CalcConc27( Distributions27 distAD; if (isEvenTimestep) { - distAD.f[dirE ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirW ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirTS *size_Mat]; + distAD.f[E ] = &distributionsAD[E *size_Mat]; + distAD.f[W ] = &distributionsAD[W *size_Mat]; + distAD.f[N ] = &distributionsAD[N *size_Mat]; + distAD.f[S ] = &distributionsAD[S *size_Mat]; + distAD.f[T ] = &distributionsAD[T *size_Mat]; + distAD.f[B ] = &distributionsAD[B *size_Mat]; + distAD.f[NE ] = &distributionsAD[NE *size_Mat]; + distAD.f[SW ] = &distributionsAD[SW *size_Mat]; + distAD.f[SE ] = &distributionsAD[SE *size_Mat]; + distAD.f[NW ] = &distributionsAD[NW *size_Mat]; + distAD.f[TE ] = &distributionsAD[TE *size_Mat]; + distAD.f[BW ] = &distributionsAD[BW *size_Mat]; + distAD.f[BE ] = &distributionsAD[BE *size_Mat]; + distAD.f[TW ] = &distributionsAD[TW *size_Mat]; + distAD.f[TN ] = &distributionsAD[TN *size_Mat]; + distAD.f[BS ] = &distributionsAD[BS *size_Mat]; + distAD.f[BN ] = &distributionsAD[BN *size_Mat]; + distAD.f[TS ] = &distributionsAD[TS *size_Mat]; distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[TNE ] = &distributionsAD[TNE *size_Mat]; + distAD.f[TSW ] = &distributionsAD[TSW *size_Mat]; + distAD.f[TSE ] = &distributionsAD[TSE *size_Mat]; + distAD.f[TNW ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNE ] = &distributionsAD[BNE *size_Mat]; + distAD.f[BSW ] = &distributionsAD[BSW *size_Mat]; + distAD.f[BSE ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNW ] = &distributionsAD[BNW *size_Mat]; } else { - distAD.f[dirW ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirE ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirTS *size_Mat]; + distAD.f[W ] = &distributionsAD[E *size_Mat]; + distAD.f[E ] = &distributionsAD[W *size_Mat]; + distAD.f[S ] = &distributionsAD[N *size_Mat]; + distAD.f[N ] = &distributionsAD[S *size_Mat]; + distAD.f[B ] = &distributionsAD[T *size_Mat]; + distAD.f[T ] = &distributionsAD[B *size_Mat]; + distAD.f[SW ] = &distributionsAD[NE *size_Mat]; + distAD.f[NE ] = &distributionsAD[SW *size_Mat]; + distAD.f[NW ] = &distributionsAD[SE *size_Mat]; + distAD.f[SE ] = &distributionsAD[NW *size_Mat]; + distAD.f[BW ] = &distributionsAD[TE *size_Mat]; + distAD.f[TE ] = &distributionsAD[BW *size_Mat]; + distAD.f[TW ] = &distributionsAD[BE *size_Mat]; + distAD.f[BE ] = &distributionsAD[TW *size_Mat]; + distAD.f[BS ] = &distributionsAD[TN *size_Mat]; + distAD.f[TN ] = &distributionsAD[BS *size_Mat]; + distAD.f[TS ] = &distributionsAD[BN *size_Mat]; + distAD.f[BN ] = &distributionsAD[TS *size_Mat]; distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirBNW *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirTSE *size_Mat]; + distAD.f[TNE ] = &distributionsAD[BSW *size_Mat]; + distAD.f[TSW ] = &distributionsAD[BNE *size_Mat]; + distAD.f[TSE ] = &distributionsAD[BNW *size_Mat]; + distAD.f[TNW ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNE ] = &distributionsAD[TSW *size_Mat]; + distAD.f[BSW ] = &distributionsAD[TNE *size_Mat]; + distAD.f[BSE ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNW ] = &distributionsAD[TSE *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) @@ -165,33 +165,33 @@ extern "C" __global__ void CalcConc27( //////////////////////////////////////////////////////////////////////////////// //! - Set local distributions //! - real mfcbb = (distAD.f[dirE ])[ke ]; - real mfabb = (distAD.f[dirW ])[kw ]; - real mfbcb = (distAD.f[dirN ])[kn ]; - real mfbab = (distAD.f[dirS ])[ks ]; - real mfbbc = (distAD.f[dirT ])[kt ]; - real mfbba = (distAD.f[dirB ])[kb ]; - real mfccb = (distAD.f[dirNE ])[kne ]; - real mfaab = (distAD.f[dirSW ])[ksw ]; - real mfcab = (distAD.f[dirSE ])[kse ]; - real mfacb = (distAD.f[dirNW ])[knw ]; - real mfcbc = (distAD.f[dirTE ])[kte ]; - real mfaba = (distAD.f[dirBW ])[kbw ]; - real mfcba = (distAD.f[dirBE ])[kbe ]; - real mfabc = (distAD.f[dirTW ])[ktw ]; - real mfbcc = (distAD.f[dirTN ])[ktn ]; - real mfbaa = (distAD.f[dirBS ])[kbs ]; - real mfbca = (distAD.f[dirBN ])[kbn ]; - real mfbac = (distAD.f[dirTS ])[kts ]; + real mfcbb = (distAD.f[E ])[ke ]; + real mfabb = (distAD.f[W ])[kw ]; + real mfbcb = (distAD.f[N ])[kn ]; + real mfbab = (distAD.f[S ])[ks ]; + real mfbbc = (distAD.f[T ])[kt ]; + real mfbba = (distAD.f[B ])[kb ]; + real mfccb = (distAD.f[NE ])[kne ]; + real mfaab = (distAD.f[SW ])[ksw ]; + real mfcab = (distAD.f[SE ])[kse ]; + real mfacb = (distAD.f[NW ])[knw ]; + real mfcbc = (distAD.f[TE ])[kte ]; + real mfaba = (distAD.f[BW ])[kbw ]; + real mfcba = (distAD.f[BE ])[kbe ]; + real mfabc = (distAD.f[TW ])[ktw ]; + real mfbcc = (distAD.f[TN ])[ktn ]; + real mfbaa = (distAD.f[BS ])[kbs ]; + real mfbca = (distAD.f[BN ])[kbn ]; + real mfbac = (distAD.f[TS ])[kts ]; real mfbbb = (distAD.f[dirREST])[k ]; - real mfccc = (distAD.f[dirTNE ])[ktne]; - real mfaac = (distAD.f[dirTSW ])[ktsw]; - real mfcac = (distAD.f[dirTSE ])[ktse]; - real mfacc = (distAD.f[dirTNW ])[ktnw]; - real mfcca = (distAD.f[dirBNE ])[kbne]; - real mfaaa = (distAD.f[dirBSW ])[kbsw]; - real mfcaa = (distAD.f[dirBSE ])[kbse]; - real mfaca = (distAD.f[dirBNW ])[kbnw]; + real mfccc = (distAD.f[TNE ])[ktne]; + real mfaac = (distAD.f[TSW ])[ktsw]; + real mfcac = (distAD.f[TSE ])[ktse]; + real mfacc = (distAD.f[TNW ])[ktnw]; + real mfcca = (distAD.f[BNE ])[kbne]; + real mfaaa = (distAD.f[BSW ])[kbsw]; + real mfcaa = (distAD.f[BSE ])[kbse]; + real mfaca = (distAD.f[BNW ])[kbnw]; ////////////////////////////////////////////////////////////////////////// //! - Calculate concentration using pyramid summation for low round-off errors as in Eq. (J1)-(J3) \ref //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa.2015.05.001 ]</b></a> diff --git a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu index 8b8e66dd6..4c590f61e 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu @@ -79,63 +79,63 @@ extern "C" __global__ void LBCalcMacCompSP27( Distributions27 dist; if (isEvenTimestep) { - dist.f[dirE ] = &distributions[dirE *size_Mat]; - dist.f[dirW ] = &distributions[dirW *size_Mat]; - dist.f[dirN ] = &distributions[dirN *size_Mat]; - dist.f[dirS ] = &distributions[dirS *size_Mat]; - dist.f[dirT ] = &distributions[dirT *size_Mat]; - dist.f[dirB ] = &distributions[dirB *size_Mat]; - dist.f[dirNE ] = &distributions[dirNE *size_Mat]; - dist.f[dirSW ] = &distributions[dirSW *size_Mat]; - dist.f[dirSE ] = &distributions[dirSE *size_Mat]; - dist.f[dirNW ] = &distributions[dirNW *size_Mat]; - dist.f[dirTE ] = &distributions[dirTE *size_Mat]; - dist.f[dirBW ] = &distributions[dirBW *size_Mat]; - dist.f[dirBE ] = &distributions[dirBE *size_Mat]; - dist.f[dirTW ] = &distributions[dirTW *size_Mat]; - dist.f[dirTN ] = &distributions[dirTN *size_Mat]; - dist.f[dirBS ] = &distributions[dirBS *size_Mat]; - dist.f[dirBN ] = &distributions[dirBN *size_Mat]; - dist.f[dirTS ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirTNE *size_Mat]; - dist.f[dirTSW ] = &distributions[dirTSW *size_Mat]; - dist.f[dirTSE ] = &distributions[dirTSE *size_Mat]; - dist.f[dirTNW ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNE ] = &distributions[dirBNE *size_Mat]; - dist.f[dirBSW ] = &distributions[dirBSW *size_Mat]; - dist.f[dirBSE ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNW ] = &distributions[dirBNW *size_Mat]; + dist.f[E ] = &distributions[E *size_Mat]; + dist.f[W ] = &distributions[W *size_Mat]; + dist.f[N ] = &distributions[N *size_Mat]; + dist.f[S ] = &distributions[S *size_Mat]; + dist.f[T ] = &distributions[T *size_Mat]; + dist.f[B ] = &distributions[B *size_Mat]; + dist.f[NE ] = &distributions[NE *size_Mat]; + dist.f[SW ] = &distributions[SW *size_Mat]; + dist.f[SE ] = &distributions[SE *size_Mat]; + dist.f[NW ] = &distributions[NW *size_Mat]; + dist.f[TE ] = &distributions[TE *size_Mat]; + dist.f[BW ] = &distributions[BW *size_Mat]; + dist.f[BE ] = &distributions[BE *size_Mat]; + dist.f[TW ] = &distributions[TW *size_Mat]; + dist.f[TN ] = &distributions[TN *size_Mat]; + dist.f[BS ] = &distributions[BS *size_Mat]; + dist.f[BN ] = &distributions[BN *size_Mat]; + dist.f[TS ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[TNE *size_Mat]; + dist.f[TSW ] = &distributions[TSW *size_Mat]; + dist.f[TSE ] = &distributions[TSE *size_Mat]; + dist.f[TNW ] = &distributions[TNW *size_Mat]; + dist.f[BNE ] = &distributions[BNE *size_Mat]; + dist.f[BSW ] = &distributions[BSW *size_Mat]; + dist.f[BSE ] = &distributions[BSE *size_Mat]; + dist.f[BNW ] = &distributions[BNW *size_Mat]; } else { - dist.f[dirW ] = &distributions[dirE *size_Mat]; - dist.f[dirE ] = &distributions[dirW *size_Mat]; - dist.f[dirS ] = &distributions[dirN *size_Mat]; - dist.f[dirN ] = &distributions[dirS *size_Mat]; - dist.f[dirB ] = &distributions[dirT *size_Mat]; - dist.f[dirT ] = &distributions[dirB *size_Mat]; - dist.f[dirSW ] = &distributions[dirNE *size_Mat]; - dist.f[dirNE ] = &distributions[dirSW *size_Mat]; - dist.f[dirNW ] = &distributions[dirSE *size_Mat]; - dist.f[dirSE ] = &distributions[dirNW *size_Mat]; - dist.f[dirBW ] = &distributions[dirTE *size_Mat]; - dist.f[dirTE ] = &distributions[dirBW *size_Mat]; - dist.f[dirTW ] = &distributions[dirBE *size_Mat]; - dist.f[dirBE ] = &distributions[dirTW *size_Mat]; - dist.f[dirBS ] = &distributions[dirTN *size_Mat]; - dist.f[dirTN ] = &distributions[dirBS *size_Mat]; - dist.f[dirTS ] = &distributions[dirBN *size_Mat]; - dist.f[dirBN ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirBSW *size_Mat]; - dist.f[dirTSW ] = &distributions[dirBNE *size_Mat]; - dist.f[dirTSE ] = &distributions[dirBNW *size_Mat]; - dist.f[dirTNW ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNE ] = &distributions[dirTSW *size_Mat]; - dist.f[dirBSW ] = &distributions[dirTNE *size_Mat]; - dist.f[dirBSE ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNW ] = &distributions[dirTSE *size_Mat]; + dist.f[W ] = &distributions[E *size_Mat]; + dist.f[E ] = &distributions[W *size_Mat]; + dist.f[S ] = &distributions[N *size_Mat]; + dist.f[N ] = &distributions[S *size_Mat]; + dist.f[B ] = &distributions[T *size_Mat]; + dist.f[T ] = &distributions[B *size_Mat]; + dist.f[SW ] = &distributions[NE *size_Mat]; + dist.f[NE ] = &distributions[SW *size_Mat]; + dist.f[NW ] = &distributions[SE *size_Mat]; + dist.f[SE ] = &distributions[NW *size_Mat]; + dist.f[BW ] = &distributions[TE *size_Mat]; + dist.f[TE ] = &distributions[BW *size_Mat]; + dist.f[TW ] = &distributions[BE *size_Mat]; + dist.f[BE ] = &distributions[TW *size_Mat]; + dist.f[BS ] = &distributions[TN *size_Mat]; + dist.f[TN ] = &distributions[BS *size_Mat]; + dist.f[TS ] = &distributions[BN *size_Mat]; + dist.f[BN ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[BSW *size_Mat]; + dist.f[TSW ] = &distributions[BNE *size_Mat]; + dist.f[TSE ] = &distributions[BNW *size_Mat]; + dist.f[TNW ] = &distributions[BSE *size_Mat]; + dist.f[BNE ] = &distributions[TSW *size_Mat]; + dist.f[BSW ] = &distributions[TNE *size_Mat]; + dist.f[BSE ] = &distributions[TNW *size_Mat]; + dist.f[BNW ] = &distributions[TSE *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) @@ -314,125 +314,125 @@ extern "C" __global__ void LBCalcMacADCompSP27( Distributions27 dist; if (isEvenTimestep) { - dist.f[dirE ] = &distributions[dirE *size_Mat]; - dist.f[dirW ] = &distributions[dirW *size_Mat]; - dist.f[dirN ] = &distributions[dirN *size_Mat]; - dist.f[dirS ] = &distributions[dirS *size_Mat]; - dist.f[dirT ] = &distributions[dirT *size_Mat]; - dist.f[dirB ] = &distributions[dirB *size_Mat]; - dist.f[dirNE ] = &distributions[dirNE *size_Mat]; - dist.f[dirSW ] = &distributions[dirSW *size_Mat]; - dist.f[dirSE ] = &distributions[dirSE *size_Mat]; - dist.f[dirNW ] = &distributions[dirNW *size_Mat]; - dist.f[dirTE ] = &distributions[dirTE *size_Mat]; - dist.f[dirBW ] = &distributions[dirBW *size_Mat]; - dist.f[dirBE ] = &distributions[dirBE *size_Mat]; - dist.f[dirTW ] = &distributions[dirTW *size_Mat]; - dist.f[dirTN ] = &distributions[dirTN *size_Mat]; - dist.f[dirBS ] = &distributions[dirBS *size_Mat]; - dist.f[dirBN ] = &distributions[dirBN *size_Mat]; - dist.f[dirTS ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirTNE *size_Mat]; - dist.f[dirTSW ] = &distributions[dirTSW *size_Mat]; - dist.f[dirTSE ] = &distributions[dirTSE *size_Mat]; - dist.f[dirTNW ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNE ] = &distributions[dirBNE *size_Mat]; - dist.f[dirBSW ] = &distributions[dirBSW *size_Mat]; - dist.f[dirBSE ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNW ] = &distributions[dirBNW *size_Mat]; + dist.f[E ] = &distributions[E *size_Mat]; + dist.f[W ] = &distributions[W *size_Mat]; + dist.f[N ] = &distributions[N *size_Mat]; + dist.f[S ] = &distributions[S *size_Mat]; + dist.f[T ] = &distributions[T *size_Mat]; + dist.f[B ] = &distributions[B *size_Mat]; + dist.f[NE ] = &distributions[NE *size_Mat]; + dist.f[SW ] = &distributions[SW *size_Mat]; + dist.f[SE ] = &distributions[SE *size_Mat]; + dist.f[NW ] = &distributions[NW *size_Mat]; + dist.f[TE ] = &distributions[TE *size_Mat]; + dist.f[BW ] = &distributions[BW *size_Mat]; + dist.f[BE ] = &distributions[BE *size_Mat]; + dist.f[TW ] = &distributions[TW *size_Mat]; + dist.f[TN ] = &distributions[TN *size_Mat]; + dist.f[BS ] = &distributions[BS *size_Mat]; + dist.f[BN ] = &distributions[BN *size_Mat]; + dist.f[TS ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[TNE *size_Mat]; + dist.f[TSW ] = &distributions[TSW *size_Mat]; + dist.f[TSE ] = &distributions[TSE *size_Mat]; + dist.f[TNW ] = &distributions[TNW *size_Mat]; + dist.f[BNE ] = &distributions[BNE *size_Mat]; + dist.f[BSW ] = &distributions[BSW *size_Mat]; + dist.f[BSE ] = &distributions[BSE *size_Mat]; + dist.f[BNW ] = &distributions[BNW *size_Mat]; } else { - dist.f[dirW ] = &distributions[dirE *size_Mat]; - dist.f[dirE ] = &distributions[dirW *size_Mat]; - dist.f[dirS ] = &distributions[dirN *size_Mat]; - dist.f[dirN ] = &distributions[dirS *size_Mat]; - dist.f[dirB ] = &distributions[dirT *size_Mat]; - dist.f[dirT ] = &distributions[dirB *size_Mat]; - dist.f[dirSW ] = &distributions[dirNE *size_Mat]; - dist.f[dirNE ] = &distributions[dirSW *size_Mat]; - dist.f[dirNW ] = &distributions[dirSE *size_Mat]; - dist.f[dirSE ] = &distributions[dirNW *size_Mat]; - dist.f[dirBW ] = &distributions[dirTE *size_Mat]; - dist.f[dirTE ] = &distributions[dirBW *size_Mat]; - dist.f[dirTW ] = &distributions[dirBE *size_Mat]; - dist.f[dirBE ] = &distributions[dirTW *size_Mat]; - dist.f[dirBS ] = &distributions[dirTN *size_Mat]; - dist.f[dirTN ] = &distributions[dirBS *size_Mat]; - dist.f[dirTS ] = &distributions[dirBN *size_Mat]; - dist.f[dirBN ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirBSW *size_Mat]; - dist.f[dirTSW ] = &distributions[dirBNE *size_Mat]; - dist.f[dirTSE ] = &distributions[dirBNW *size_Mat]; - dist.f[dirTNW ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNE ] = &distributions[dirTSW *size_Mat]; - dist.f[dirBSW ] = &distributions[dirTNE *size_Mat]; - dist.f[dirBSE ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNW ] = &distributions[dirTSE *size_Mat]; + dist.f[W ] = &distributions[E *size_Mat]; + dist.f[E ] = &distributions[W *size_Mat]; + dist.f[S ] = &distributions[N *size_Mat]; + dist.f[N ] = &distributions[S *size_Mat]; + dist.f[B ] = &distributions[T *size_Mat]; + dist.f[T ] = &distributions[B *size_Mat]; + dist.f[SW ] = &distributions[NE *size_Mat]; + dist.f[NE ] = &distributions[SW *size_Mat]; + dist.f[NW ] = &distributions[SE *size_Mat]; + dist.f[SE ] = &distributions[NW *size_Mat]; + dist.f[BW ] = &distributions[TE *size_Mat]; + dist.f[TE ] = &distributions[BW *size_Mat]; + dist.f[TW ] = &distributions[BE *size_Mat]; + dist.f[BE ] = &distributions[TW *size_Mat]; + dist.f[BS ] = &distributions[TN *size_Mat]; + dist.f[TN ] = &distributions[BS *size_Mat]; + dist.f[TS ] = &distributions[BN *size_Mat]; + dist.f[BN ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[BSW *size_Mat]; + dist.f[TSW ] = &distributions[BNE *size_Mat]; + dist.f[TSE ] = &distributions[BNW *size_Mat]; + dist.f[TNW ] = &distributions[BSE *size_Mat]; + dist.f[BNE ] = &distributions[TSW *size_Mat]; + dist.f[BSW ] = &distributions[TNE *size_Mat]; + dist.f[BSE ] = &distributions[TNW *size_Mat]; + dist.f[BNW ] = &distributions[TSE *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// Distributions27 distAD; if (isEvenTimestep) { - distAD.f[dirE ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirW ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[E ] = &distributionsAD[E *size_Mat]; + distAD.f[W ] = &distributionsAD[W *size_Mat]; + distAD.f[N ] = &distributionsAD[N *size_Mat]; + distAD.f[S ] = &distributionsAD[S *size_Mat]; + distAD.f[T ] = &distributionsAD[T *size_Mat]; + distAD.f[B ] = &distributionsAD[B *size_Mat]; + distAD.f[NE ] = &distributionsAD[NE *size_Mat]; + distAD.f[SW ] = &distributionsAD[SW *size_Mat]; + distAD.f[SE ] = &distributionsAD[SE *size_Mat]; + distAD.f[NW ] = &distributionsAD[NW *size_Mat]; + distAD.f[TE ] = &distributionsAD[TE *size_Mat]; + distAD.f[BW ] = &distributionsAD[BW *size_Mat]; + distAD.f[BE ] = &distributionsAD[BE *size_Mat]; + distAD.f[TW ] = &distributionsAD[TW *size_Mat]; + distAD.f[TN ] = &distributionsAD[TN *size_Mat]; + distAD.f[BS ] = &distributionsAD[BS *size_Mat]; + distAD.f[BN ] = &distributionsAD[BN *size_Mat]; + distAD.f[TS ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[TNE ] = &distributionsAD[TNE *size_Mat]; + distAD.f[TSW ] = &distributionsAD[TSW *size_Mat]; + distAD.f[TSE ] = &distributionsAD[TSE *size_Mat]; + distAD.f[TNW ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNE ] = &distributionsAD[BNE *size_Mat]; + distAD.f[BSW ] = &distributionsAD[BSW *size_Mat]; + distAD.f[BSE ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNW ] = &distributionsAD[BNW *size_Mat]; } else { - distAD.f[dirW ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirE ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[W ] = &distributionsAD[E *size_Mat]; + distAD.f[E ] = &distributionsAD[W *size_Mat]; + distAD.f[S ] = &distributionsAD[N *size_Mat]; + distAD.f[N ] = &distributionsAD[S *size_Mat]; + distAD.f[B ] = &distributionsAD[T *size_Mat]; + distAD.f[T ] = &distributionsAD[B *size_Mat]; + distAD.f[SW ] = &distributionsAD[NE *size_Mat]; + distAD.f[NE ] = &distributionsAD[SW *size_Mat]; + distAD.f[NW ] = &distributionsAD[SE *size_Mat]; + distAD.f[SE ] = &distributionsAD[NW *size_Mat]; + distAD.f[BW ] = &distributionsAD[TE *size_Mat]; + distAD.f[TE ] = &distributionsAD[BW *size_Mat]; + distAD.f[TW ] = &distributionsAD[BE *size_Mat]; + distAD.f[BE ] = &distributionsAD[TW *size_Mat]; + distAD.f[BS ] = &distributionsAD[TN *size_Mat]; + distAD.f[TN ] = &distributionsAD[BS *size_Mat]; + distAD.f[TS ] = &distributionsAD[BN *size_Mat]; + distAD.f[BN ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[BSW ] = &distributionsAD[TNE *size_Mat]; + distAD.f[BNE ] = &distributionsAD[TSW *size_Mat]; + distAD.f[BNW ] = &distributionsAD[TSE *size_Mat]; + distAD.f[BSE ] = &distributionsAD[TNW *size_Mat]; + distAD.f[TSW ] = &distributionsAD[BNE *size_Mat]; + distAD.f[TNE ] = &distributionsAD[BSW *size_Mat]; + distAD.f[TNW ] = &distributionsAD[BSE *size_Mat]; + distAD.f[TSE ] = &distributionsAD[BNW *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) diff --git a/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu b/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu index a08da5ba3..a767d2902 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/Cumulant27chim.cu @@ -138,63 +138,63 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel( Distributions27 dist; if (isEvenTimestep) { - dist.f[dirE ] = &distributions[dirE *size_Mat]; - dist.f[dirW ] = &distributions[dirW *size_Mat]; - dist.f[dirN ] = &distributions[dirN *size_Mat]; - dist.f[dirS ] = &distributions[dirS *size_Mat]; - dist.f[dirT ] = &distributions[dirT *size_Mat]; - dist.f[dirB ] = &distributions[dirB *size_Mat]; - dist.f[dirNE ] = &distributions[dirNE *size_Mat]; - dist.f[dirSW ] = &distributions[dirSW *size_Mat]; - dist.f[dirSE ] = &distributions[dirSE *size_Mat]; - dist.f[dirNW ] = &distributions[dirNW *size_Mat]; - dist.f[dirTE ] = &distributions[dirTE *size_Mat]; - dist.f[dirBW ] = &distributions[dirBW *size_Mat]; - dist.f[dirBE ] = &distributions[dirBE *size_Mat]; - dist.f[dirTW ] = &distributions[dirTW *size_Mat]; - dist.f[dirTN ] = &distributions[dirTN *size_Mat]; - dist.f[dirBS ] = &distributions[dirBS *size_Mat]; - dist.f[dirBN ] = &distributions[dirBN *size_Mat]; - dist.f[dirTS ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirTNE *size_Mat]; - dist.f[dirTSW ] = &distributions[dirTSW *size_Mat]; - dist.f[dirTSE ] = &distributions[dirTSE *size_Mat]; - dist.f[dirTNW ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNE ] = &distributions[dirBNE *size_Mat]; - dist.f[dirBSW ] = &distributions[dirBSW *size_Mat]; - dist.f[dirBSE ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNW ] = &distributions[dirBNW *size_Mat]; + dist.f[E ] = &distributions[E *size_Mat]; + dist.f[W ] = &distributions[W *size_Mat]; + dist.f[N ] = &distributions[N *size_Mat]; + dist.f[S ] = &distributions[S *size_Mat]; + dist.f[T ] = &distributions[T *size_Mat]; + dist.f[B ] = &distributions[B *size_Mat]; + dist.f[NE ] = &distributions[NE *size_Mat]; + dist.f[SW ] = &distributions[SW *size_Mat]; + dist.f[SE ] = &distributions[SE *size_Mat]; + dist.f[NW ] = &distributions[NW *size_Mat]; + dist.f[TE ] = &distributions[TE *size_Mat]; + dist.f[BW ] = &distributions[BW *size_Mat]; + dist.f[BE ] = &distributions[BE *size_Mat]; + dist.f[TW ] = &distributions[TW *size_Mat]; + dist.f[TN ] = &distributions[TN *size_Mat]; + dist.f[BS ] = &distributions[BS *size_Mat]; + dist.f[BN ] = &distributions[BN *size_Mat]; + dist.f[TS ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[TNE *size_Mat]; + dist.f[TSW ] = &distributions[TSW *size_Mat]; + dist.f[TSE ] = &distributions[TSE *size_Mat]; + dist.f[TNW ] = &distributions[TNW *size_Mat]; + dist.f[BNE ] = &distributions[BNE *size_Mat]; + dist.f[BSW ] = &distributions[BSW *size_Mat]; + dist.f[BSE ] = &distributions[BSE *size_Mat]; + dist.f[BNW ] = &distributions[BNW *size_Mat]; } else { - dist.f[dirW ] = &distributions[dirE *size_Mat]; - dist.f[dirE ] = &distributions[dirW *size_Mat]; - dist.f[dirS ] = &distributions[dirN *size_Mat]; - dist.f[dirN ] = &distributions[dirS *size_Mat]; - dist.f[dirB ] = &distributions[dirT *size_Mat]; - dist.f[dirT ] = &distributions[dirB *size_Mat]; - dist.f[dirSW ] = &distributions[dirNE *size_Mat]; - dist.f[dirNE ] = &distributions[dirSW *size_Mat]; - dist.f[dirNW ] = &distributions[dirSE *size_Mat]; - dist.f[dirSE ] = &distributions[dirNW *size_Mat]; - dist.f[dirBW ] = &distributions[dirTE *size_Mat]; - dist.f[dirTE ] = &distributions[dirBW *size_Mat]; - dist.f[dirTW ] = &distributions[dirBE *size_Mat]; - dist.f[dirBE ] = &distributions[dirTW *size_Mat]; - dist.f[dirBS ] = &distributions[dirTN *size_Mat]; - dist.f[dirTN ] = &distributions[dirBS *size_Mat]; - dist.f[dirTS ] = &distributions[dirBN *size_Mat]; - dist.f[dirBN ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirBSW ] = &distributions[dirTNE *size_Mat]; - dist.f[dirBNE ] = &distributions[dirTSW *size_Mat]; - dist.f[dirBNW ] = &distributions[dirTSE *size_Mat]; - dist.f[dirBSE ] = &distributions[dirTNW *size_Mat]; - dist.f[dirTSW ] = &distributions[dirBNE *size_Mat]; - dist.f[dirTNE ] = &distributions[dirBSW *size_Mat]; - dist.f[dirTNW ] = &distributions[dirBSE *size_Mat]; - dist.f[dirTSE ] = &distributions[dirBNW *size_Mat]; + dist.f[W ] = &distributions[E *size_Mat]; + dist.f[E ] = &distributions[W *size_Mat]; + dist.f[S ] = &distributions[N *size_Mat]; + dist.f[N ] = &distributions[S *size_Mat]; + dist.f[B ] = &distributions[T *size_Mat]; + dist.f[T ] = &distributions[B *size_Mat]; + dist.f[SW ] = &distributions[NE *size_Mat]; + dist.f[NE ] = &distributions[SW *size_Mat]; + dist.f[NW ] = &distributions[SE *size_Mat]; + dist.f[SE ] = &distributions[NW *size_Mat]; + dist.f[BW ] = &distributions[TE *size_Mat]; + dist.f[TE ] = &distributions[BW *size_Mat]; + dist.f[TW ] = &distributions[BE *size_Mat]; + dist.f[BE ] = &distributions[TW *size_Mat]; + dist.f[BS ] = &distributions[TN *size_Mat]; + dist.f[TN ] = &distributions[BS *size_Mat]; + dist.f[TS ] = &distributions[BN *size_Mat]; + dist.f[BN ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[BSW ] = &distributions[TNE *size_Mat]; + dist.f[BNE ] = &distributions[TSW *size_Mat]; + dist.f[BNW ] = &distributions[TSE *size_Mat]; + dist.f[BSE ] = &distributions[TNW *size_Mat]; + dist.f[TSW ] = &distributions[BNE *size_Mat]; + dist.f[TNE ] = &distributions[BSW *size_Mat]; + dist.f[TNW ] = &distributions[BSE *size_Mat]; + dist.f[TSE ] = &distributions[BNW *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) @@ -208,33 +208,33 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel( //////////////////////////////////////////////////////////////////////////////////// //! - Set local distributions //! - real mfcbb = (dist.f[dirE ])[k]; - real mfabb = (dist.f[dirW ])[kw]; - real mfbcb = (dist.f[dirN ])[k]; - real mfbab = (dist.f[dirS ])[ks]; - real mfbbc = (dist.f[dirT ])[k]; - real mfbba = (dist.f[dirB ])[kb]; - real mfccb = (dist.f[dirNE ])[k]; - real mfaab = (dist.f[dirSW ])[ksw]; - real mfcab = (dist.f[dirSE ])[ks]; - real mfacb = (dist.f[dirNW ])[kw]; - real mfcbc = (dist.f[dirTE ])[k]; - real mfaba = (dist.f[dirBW ])[kbw]; - real mfcba = (dist.f[dirBE ])[kb]; - real mfabc = (dist.f[dirTW ])[kw]; - real mfbcc = (dist.f[dirTN ])[k]; - real mfbaa = (dist.f[dirBS ])[kbs]; - real mfbca = (dist.f[dirBN ])[kb]; - real mfbac = (dist.f[dirTS ])[ks]; - real mfbbb = (dist.f[dirREST])[k]; - real mfccc = (dist.f[dirTNE ])[k]; - real mfaac = (dist.f[dirTSW ])[ksw]; - real mfcac = (dist.f[dirTSE ])[ks]; - real mfacc = (dist.f[dirTNW ])[kw]; - real mfcca = (dist.f[dirBNE ])[kb]; - real mfaaa = (dist.f[dirBSW ])[kbsw]; - real mfcaa = (dist.f[dirBSE ])[kbs]; - real mfaca = (dist.f[dirBNW ])[kbw]; + real mfcbb = (dist.f[E ])[k]; + real mfabb = (dist.f[W ])[kw]; + real mfbcb = (dist.f[N ])[k]; + real mfbab = (dist.f[S ])[ks]; + real mfbbc = (dist.f[T ])[k]; + real mfbba = (dist.f[B ])[kb]; + real mfccb = (dist.f[NE ])[k]; + real mfaab = (dist.f[SW ])[ksw]; + real mfcab = (dist.f[SE ])[ks]; + real mfacb = (dist.f[NW ])[kw]; + real mfcbc = (dist.f[TE ])[k]; + real mfaba = (dist.f[BW ])[kbw]; + real mfcba = (dist.f[BE ])[kb]; + real mfabc = (dist.f[TW ])[kw]; + real mfbcc = (dist.f[TN ])[k]; + real mfbaa = (dist.f[BS ])[kbs]; + real mfbca = (dist.f[BN ])[kb]; + real mfbac = (dist.f[TS ])[ks]; + real mfbbb = (dist.f[REST])[k]; + real mfccc = (dist.f[TNE ])[k]; + real mfaac = (dist.f[TSW ])[ksw]; + real mfcac = (dist.f[TSE ])[ks]; + real mfacc = (dist.f[TNW ])[kw]; + real mfcca = (dist.f[BNE ])[kb]; + real mfaaa = (dist.f[BSW ])[kbsw]; + real mfcaa = (dist.f[BSE ])[kbs]; + real mfaca = (dist.f[BNW ])[kbw]; //////////////////////////////////////////////////////////////////////////////////// //! - Calculate density and velocity using pyramid summation for low round-off errors as in Eq. (J1)-(J3) \ref //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), DOI:10.1016/j.camwa.2015.05.001 ]</b></a> @@ -610,33 +610,33 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel( //! stored arrays dependent on timestep is based on the esoteric twist algorithm //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), DOI:10.3390/computation5020019 ]</b></a> //! - (dist.f[dirE ])[k ] = mfabb; - (dist.f[dirW ])[kw ] = mfcbb; - (dist.f[dirN ])[k ] = mfbab; - (dist.f[dirS ])[ks ] = mfbcb; - (dist.f[dirT ])[k ] = mfbba; - (dist.f[dirB ])[kb ] = mfbbc; - (dist.f[dirNE ])[k ] = mfaab; - (dist.f[dirSW ])[ksw ] = mfccb; - (dist.f[dirSE ])[ks ] = mfacb; - (dist.f[dirNW ])[kw ] = mfcab; - (dist.f[dirTE ])[k ] = mfaba; - (dist.f[dirBW ])[kbw ] = mfcbc; - (dist.f[dirBE ])[kb ] = mfabc; - (dist.f[dirTW ])[kw ] = mfcba; - (dist.f[dirTN ])[k ] = mfbaa; - (dist.f[dirBS ])[kbs ] = mfbcc; - (dist.f[dirBN ])[kb ] = mfbac; - (dist.f[dirTS ])[ks ] = mfbca; - (dist.f[dirREST])[k ] = mfbbb; - (dist.f[dirTNE ])[k ] = mfaaa; - (dist.f[dirTSE ])[ks ] = mfaca; - (dist.f[dirBNE ])[kb ] = mfaac; - (dist.f[dirBSE ])[kbs ] = mfacc; - (dist.f[dirTNW ])[kw ] = mfcaa; - (dist.f[dirTSW ])[ksw ] = mfcca; - (dist.f[dirBNW ])[kbw ] = mfcac; - (dist.f[dirBSW ])[kbsw] = mfccc; + (dist.f[E ])[k ] = mfabb; + (dist.f[W ])[kw ] = mfcbb; + (dist.f[N ])[k ] = mfbab; + (dist.f[S ])[ks ] = mfbcb; + (dist.f[T ])[k ] = mfbba; + (dist.f[B ])[kb ] = mfbbc; + (dist.f[NE ])[k ] = mfaab; + (dist.f[SW ])[ksw ] = mfccb; + (dist.f[SE ])[ks ] = mfacb; + (dist.f[NW ])[kw ] = mfcab; + (dist.f[TE ])[k ] = mfaba; + (dist.f[BW ])[kbw ] = mfcbc; + (dist.f[BE ])[kb ] = mfabc; + (dist.f[TW ])[kw ] = mfcba; + (dist.f[TN ])[k ] = mfbaa; + (dist.f[BS ])[kbs ] = mfbcc; + (dist.f[BN ])[kb ] = mfbac; + (dist.f[TS ])[ks ] = mfbca; + (dist.f[REST])[k ] = mfbbb; + (dist.f[TNE ])[k ] = mfaaa; + (dist.f[TSE ])[ks ] = mfaca; + (dist.f[BNE ])[kb ] = mfaac; + (dist.f[BSE ])[kbs ] = mfacc; + (dist.f[TNW ])[kw ] = mfcaa; + (dist.f[TSW ])[ksw ] = mfcca; + (dist.f[BNW ])[kbw ] = mfcac; + (dist.f[BSW ])[kbsw] = mfccc; } } //////////////////////////////////////////////////////////////////////////////// @@ -771,63 +771,63 @@ extern "C" __global__ void Cumulant_K17_LBM_Device_Kernel_AD( Distributions27 distAD; if (isEvenTimestep) { - distAD.f[dirE ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirW ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[E ] = &distributionsAD[E *size_Mat]; + distAD.f[W ] = &distributionsAD[W *size_Mat]; + distAD.f[N ] = &distributionsAD[N *size_Mat]; + distAD.f[S ] = &distributionsAD[S *size_Mat]; + distAD.f[T ] = &distributionsAD[T *size_Mat]; + distAD.f[B ] = &distributionsAD[B *size_Mat]; + distAD.f[NE ] = &distributionsAD[NE *size_Mat]; + distAD.f[SW ] = &distributionsAD[SW *size_Mat]; + distAD.f[SE ] = &distributionsAD[SE *size_Mat]; + distAD.f[NW ] = &distributionsAD[NW *size_Mat]; + distAD.f[TE ] = &distributionsAD[TE *size_Mat]; + distAD.f[BW ] = &distributionsAD[BW *size_Mat]; + distAD.f[BE ] = &distributionsAD[BE *size_Mat]; + distAD.f[TW ] = &distributionsAD[TW *size_Mat]; + distAD.f[TN ] = &distributionsAD[TN *size_Mat]; + distAD.f[BS ] = &distributionsAD[BS *size_Mat]; + distAD.f[BN ] = &distributionsAD[BN *size_Mat]; + distAD.f[TS ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[TNE ] = &distributionsAD[TNE *size_Mat]; + distAD.f[TSW ] = &distributionsAD[TSW *size_Mat]; + distAD.f[TSE ] = &distributionsAD[TSE *size_Mat]; + distAD.f[TNW ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNE ] = &distributionsAD[BNE *size_Mat]; + distAD.f[BSW ] = &distributionsAD[BSW *size_Mat]; + distAD.f[BSE ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNW ] = &distributionsAD[BNW *size_Mat]; } else { - distAD.f[dirW ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirE ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[W ] = &distributionsAD[E *size_Mat]; + distAD.f[E ] = &distributionsAD[W *size_Mat]; + distAD.f[S ] = &distributionsAD[N *size_Mat]; + distAD.f[N ] = &distributionsAD[S *size_Mat]; + distAD.f[B ] = &distributionsAD[T *size_Mat]; + distAD.f[T ] = &distributionsAD[B *size_Mat]; + distAD.f[SW ] = &distributionsAD[NE *size_Mat]; + distAD.f[NE ] = &distributionsAD[SW *size_Mat]; + distAD.f[NW ] = &distributionsAD[SE *size_Mat]; + distAD.f[SE ] = &distributionsAD[NW *size_Mat]; + distAD.f[BW ] = &distributionsAD[TE *size_Mat]; + distAD.f[TE ] = &distributionsAD[BW *size_Mat]; + distAD.f[TW ] = &distributionsAD[BE *size_Mat]; + distAD.f[BE ] = &distributionsAD[TW *size_Mat]; + distAD.f[BS ] = &distributionsAD[TN *size_Mat]; + distAD.f[TN ] = &distributionsAD[BS *size_Mat]; + distAD.f[TS ] = &distributionsAD[BN *size_Mat]; + distAD.f[BN ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[BSW ] = &distributionsAD[TNE *size_Mat]; + distAD.f[BNE ] = &distributionsAD[TSW *size_Mat]; + distAD.f[BNW ] = &distributionsAD[TSE *size_Mat]; + distAD.f[BSE ] = &distributionsAD[TNW *size_Mat]; + distAD.f[TSW ] = &distributionsAD[BNE *size_Mat]; + distAD.f[TNE ] = &distributionsAD[BSW *size_Mat]; + distAD.f[TNW ] = &distributionsAD[BSE *size_Mat]; + distAD.f[TSE ] = &distributionsAD[BNW *size_Mat]; } //////////////////////////////////////////////////////////////////////////////// //! - Set neighbor indices (necessary for indirect addressing) diff --git a/src/gpu/VirtualFluids_GPU/GPU/Init27.cu b/src/gpu/VirtualFluids_GPU/GPU/Init27.cu index 08660a5c5..2d0a76ef3 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/Init27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/Init27.cu @@ -76,63 +76,63 @@ extern "C" __global__ void LBInit( Distributions27 dist; if (isEvenTimestep) { - dist.f[dirE ] = &distributions[dirE *size_Mat]; - dist.f[dirW ] = &distributions[dirW *size_Mat]; - dist.f[dirN ] = &distributions[dirN *size_Mat]; - dist.f[dirS ] = &distributions[dirS *size_Mat]; - dist.f[dirT ] = &distributions[dirT *size_Mat]; - dist.f[dirB ] = &distributions[dirB *size_Mat]; - dist.f[dirNE ] = &distributions[dirNE *size_Mat]; - dist.f[dirSW ] = &distributions[dirSW *size_Mat]; - dist.f[dirSE ] = &distributions[dirSE *size_Mat]; - dist.f[dirNW ] = &distributions[dirNW *size_Mat]; - dist.f[dirTE ] = &distributions[dirTE *size_Mat]; - dist.f[dirBW ] = &distributions[dirBW *size_Mat]; - dist.f[dirBE ] = &distributions[dirBE *size_Mat]; - dist.f[dirTW ] = &distributions[dirTW *size_Mat]; - dist.f[dirTN ] = &distributions[dirTN *size_Mat]; - dist.f[dirBS ] = &distributions[dirBS *size_Mat]; - dist.f[dirBN ] = &distributions[dirBN *size_Mat]; - dist.f[dirTS ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirTNE ] = &distributions[dirTNE *size_Mat]; - dist.f[dirTSW ] = &distributions[dirTSW *size_Mat]; - dist.f[dirTSE ] = &distributions[dirTSE *size_Mat]; - dist.f[dirTNW ] = &distributions[dirTNW *size_Mat]; - dist.f[dirBNE ] = &distributions[dirBNE *size_Mat]; - dist.f[dirBSW ] = &distributions[dirBSW *size_Mat]; - dist.f[dirBSE ] = &distributions[dirBSE *size_Mat]; - dist.f[dirBNW ] = &distributions[dirBNW *size_Mat]; + dist.f[E ] = &distributions[E *size_Mat]; + dist.f[W ] = &distributions[W *size_Mat]; + dist.f[N ] = &distributions[N *size_Mat]; + dist.f[S ] = &distributions[S *size_Mat]; + dist.f[T ] = &distributions[T *size_Mat]; + dist.f[B ] = &distributions[B *size_Mat]; + dist.f[NE ] = &distributions[NE *size_Mat]; + dist.f[SW ] = &distributions[SW *size_Mat]; + dist.f[SE ] = &distributions[SE *size_Mat]; + dist.f[NW ] = &distributions[NW *size_Mat]; + dist.f[TE ] = &distributions[TE *size_Mat]; + dist.f[BW ] = &distributions[BW *size_Mat]; + dist.f[BE ] = &distributions[BE *size_Mat]; + dist.f[TW ] = &distributions[TW *size_Mat]; + dist.f[TN ] = &distributions[TN *size_Mat]; + dist.f[BS ] = &distributions[BS *size_Mat]; + dist.f[BN ] = &distributions[BN *size_Mat]; + dist.f[TS ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[TNE ] = &distributions[TNE *size_Mat]; + dist.f[TSW ] = &distributions[TSW *size_Mat]; + dist.f[TSE ] = &distributions[TSE *size_Mat]; + dist.f[TNW ] = &distributions[TNW *size_Mat]; + dist.f[BNE ] = &distributions[BNE *size_Mat]; + dist.f[BSW ] = &distributions[BSW *size_Mat]; + dist.f[BSE ] = &distributions[BSE *size_Mat]; + dist.f[BNW ] = &distributions[BNW *size_Mat]; } else { - dist.f[dirW ] = &distributions[dirE *size_Mat]; - dist.f[dirE ] = &distributions[dirW *size_Mat]; - dist.f[dirS ] = &distributions[dirN *size_Mat]; - dist.f[dirN ] = &distributions[dirS *size_Mat]; - dist.f[dirB ] = &distributions[dirT *size_Mat]; - dist.f[dirT ] = &distributions[dirB *size_Mat]; - dist.f[dirSW ] = &distributions[dirNE *size_Mat]; - dist.f[dirNE ] = &distributions[dirSW *size_Mat]; - dist.f[dirNW ] = &distributions[dirSE *size_Mat]; - dist.f[dirSE ] = &distributions[dirNW *size_Mat]; - dist.f[dirBW ] = &distributions[dirTE *size_Mat]; - dist.f[dirTE ] = &distributions[dirBW *size_Mat]; - dist.f[dirTW ] = &distributions[dirBE *size_Mat]; - dist.f[dirBE ] = &distributions[dirTW *size_Mat]; - dist.f[dirBS ] = &distributions[dirTN *size_Mat]; - dist.f[dirTN ] = &distributions[dirBS *size_Mat]; - dist.f[dirTS ] = &distributions[dirBN *size_Mat]; - dist.f[dirBN ] = &distributions[dirTS *size_Mat]; - dist.f[dirREST] = &distributions[dirREST*size_Mat]; - dist.f[dirBSW ] = &distributions[dirTNE *size_Mat]; - dist.f[dirBNE ] = &distributions[dirTSW *size_Mat]; - dist.f[dirBNW ] = &distributions[dirTSE *size_Mat]; - dist.f[dirBSE ] = &distributions[dirTNW *size_Mat]; - dist.f[dirTSW ] = &distributions[dirBNE *size_Mat]; - dist.f[dirTNE ] = &distributions[dirBSW *size_Mat]; - dist.f[dirTNW ] = &distributions[dirBSE *size_Mat]; - dist.f[dirTSE ] = &distributions[dirBNW *size_Mat]; + dist.f[W ] = &distributions[E *size_Mat]; + dist.f[E ] = &distributions[W *size_Mat]; + dist.f[S ] = &distributions[N *size_Mat]; + dist.f[N ] = &distributions[S *size_Mat]; + dist.f[B ] = &distributions[T *size_Mat]; + dist.f[T ] = &distributions[B *size_Mat]; + dist.f[SW ] = &distributions[NE *size_Mat]; + dist.f[NE ] = &distributions[SW *size_Mat]; + dist.f[NW ] = &distributions[SE *size_Mat]; + dist.f[SE ] = &distributions[NW *size_Mat]; + dist.f[BW ] = &distributions[TE *size_Mat]; + dist.f[TE ] = &distributions[BW *size_Mat]; + dist.f[TW ] = &distributions[BE *size_Mat]; + dist.f[BE ] = &distributions[TW *size_Mat]; + dist.f[BS ] = &distributions[TN *size_Mat]; + dist.f[TN ] = &distributions[BS *size_Mat]; + dist.f[TS ] = &distributions[BN *size_Mat]; + dist.f[BN ] = &distributions[TS *size_Mat]; + dist.f[REST] = &distributions[REST*size_Mat]; + dist.f[BSW ] = &distributions[TNE *size_Mat]; + dist.f[BNE ] = &distributions[TSW *size_Mat]; + dist.f[BNW ] = &distributions[TSE *size_Mat]; + dist.f[BSE ] = &distributions[TNW *size_Mat]; + dist.f[TSW ] = &distributions[BNE *size_Mat]; + dist.f[TNE ] = &distributions[BSW *size_Mat]; + dist.f[TNW ] = &distributions[BSE *size_Mat]; + dist.f[TSE ] = &distributions[BNW *size_Mat]; } ////////////////////////////////////////////////////////////////////////// //! - Set local velocities and density @@ -176,33 +176,33 @@ extern "C" __global__ void LBInit( //! real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3); - (dist.f[dirREST])[kzero] = c8o27* (drho - cu_sq); - (dist.f[dirE ])[ke ] = c2o27* (drho + c3o1*( vx1 ) + c9o2*( vx1 )*( vx1 ) - cu_sq); - (dist.f[dirW ])[kw ] = c2o27* (drho + c3o1*(-vx1 ) + c9o2*(-vx1 )*(-vx1 ) - cu_sq); - (dist.f[dirN ])[kn ] = c2o27* (drho + c3o1*( vx2 ) + c9o2*( vx2 )*( vx2 ) - cu_sq); - (dist.f[dirS ])[ks ] = c2o27* (drho + c3o1*( - vx2 ) + c9o2*( - vx2 )*( - vx2 ) - cu_sq); - (dist.f[dirT ])[kt ] = c2o27* (drho + c3o1*( vx3) + c9o2*( vx3)*( vx3) - cu_sq); - (dist.f[dirB ])[kb ] = c2o27* (drho + c3o1*( -vx3) + c9o2*( -vx3)*( -vx3) - cu_sq); - (dist.f[dirNE ])[kne ] = c1o54* (drho + c3o1*( vx1 + vx2 ) + c9o2*( vx1 + vx2 )*( vx1 + vx2 ) - cu_sq); - (dist.f[dirSW ])[ksw ] = c1o54* (drho + c3o1*(-vx1 - vx2 ) + c9o2*(-vx1 - vx2 )*(-vx1 - vx2 ) - cu_sq); - (dist.f[dirSE ])[kse ] = c1o54* (drho + c3o1*( vx1 - vx2 ) + c9o2*( vx1 - vx2 )*( vx1 - vx2 ) - cu_sq); - (dist.f[dirNW ])[knw ] = c1o54* (drho + c3o1*(-vx1 + vx2 ) + c9o2*(-vx1 + vx2 )*(-vx1 + vx2 ) - cu_sq); - (dist.f[dirTE ])[kte ] = c1o54* (drho + c3o1*( vx1 + vx3) + c9o2*( vx1 + vx3)*( vx1 + vx3) - cu_sq); - (dist.f[dirBW ])[kbw ] = c1o54* (drho + c3o1*(-vx1 - vx3) + c9o2*(-vx1 - vx3)*(-vx1 - vx3) - cu_sq); - (dist.f[dirBE ])[kbe ] = c1o54* (drho + c3o1*( vx1 - vx3) + c9o2*( vx1 - vx3)*( vx1 - vx3) - cu_sq); - (dist.f[dirTW ])[ktw ] = c1o54* (drho + c3o1*(-vx1 + vx3) + c9o2*(-vx1 + vx3)*(-vx1 + vx3) - cu_sq); - (dist.f[dirTN ])[ktn ] = c1o54* (drho + c3o1*( vx2 + vx3) + c9o2*( vx2 + vx3)*( vx2 + vx3) - cu_sq); - (dist.f[dirBS ])[kbs ] = c1o54* (drho + c3o1*( - vx2 - vx3) + c9o2*( - vx2 - vx3)*( - vx2 - vx3) - cu_sq); - (dist.f[dirBN ])[kbn ] = c1o54* (drho + c3o1*( vx2 - vx3) + c9o2*( vx2 - vx3)*( vx2 - vx3) - cu_sq); - (dist.f[dirTS ])[kts ] = c1o54* (drho + c3o1*( - vx2 + vx3) + c9o2*( - vx2 + vx3)*( - vx2 + vx3) - cu_sq); - (dist.f[dirTNE ])[ktne ] = c1o216*(drho + c3o1*( vx1 + vx2 + vx3) + c9o2*( vx1 + vx2 + vx3)*( vx1 + vx2 + vx3) - cu_sq); - (dist.f[dirBSW ])[kbsw ] = c1o216*(drho + c3o1*(-vx1 - vx2 - vx3) + c9o2*(-vx1 - vx2 - vx3)*(-vx1 - vx2 - vx3) - cu_sq); - (dist.f[dirBNE ])[kbne ] = c1o216*(drho + c3o1*( vx1 + vx2 - vx3) + c9o2*( vx1 + vx2 - vx3)*( vx1 + vx2 - vx3) - cu_sq); - (dist.f[dirTSW ])[ktsw ] = c1o216*(drho + c3o1*(-vx1 - vx2 + vx3) + c9o2*(-vx1 - vx2 + vx3)*(-vx1 - vx2 + vx3) - cu_sq); - (dist.f[dirTSE ])[ktse ] = c1o216*(drho + c3o1*( vx1 - vx2 + vx3) + c9o2*( vx1 - vx2 + vx3)*( vx1 - vx2 + vx3) - cu_sq); - (dist.f[dirBNW ])[kbnw ] = c1o216*(drho + c3o1*(-vx1 + vx2 - vx3) + c9o2*(-vx1 + vx2 - vx3)*(-vx1 + vx2 - vx3) - cu_sq); - (dist.f[dirBSE ])[kbse ] = c1o216*(drho + c3o1*( vx1 - vx2 - vx3) + c9o2*( vx1 - vx2 - vx3)*( vx1 - vx2 - vx3) - cu_sq); - (dist.f[dirTNW ])[ktnw ] = c1o216*(drho + c3o1*(-vx1 + vx2 + vx3) + c9o2*(-vx1 + vx2 + vx3)*(-vx1 + vx2 + vx3) - cu_sq); + (dist.f[REST])[kzero] = c8o27* (drho - cu_sq); + (dist.f[E ])[ke ] = c2o27* (drho + c3o1*( vx1 ) + c9o2*( vx1 )*( vx1 ) - cu_sq); + (dist.f[W ])[kw ] = c2o27* (drho + c3o1*(-vx1 ) + c9o2*(-vx1 )*(-vx1 ) - cu_sq); + (dist.f[N ])[kn ] = c2o27* (drho + c3o1*( vx2 ) + c9o2*( vx2 )*( vx2 ) - cu_sq); + (dist.f[S ])[ks ] = c2o27* (drho + c3o1*( - vx2 ) + c9o2*( - vx2 )*( - vx2 ) - cu_sq); + (dist.f[T ])[kt ] = c2o27* (drho + c3o1*( vx3) + c9o2*( vx3)*( vx3) - cu_sq); + (dist.f[B ])[kb ] = c2o27* (drho + c3o1*( -vx3) + c9o2*( -vx3)*( -vx3) - cu_sq); + (dist.f[NE ])[kne ] = c1o54* (drho + c3o1*( vx1 + vx2 ) + c9o2*( vx1 + vx2 )*( vx1 + vx2 ) - cu_sq); + (dist.f[SW ])[ksw ] = c1o54* (drho + c3o1*(-vx1 - vx2 ) + c9o2*(-vx1 - vx2 )*(-vx1 - vx2 ) - cu_sq); + (dist.f[SE ])[kse ] = c1o54* (drho + c3o1*( vx1 - vx2 ) + c9o2*( vx1 - vx2 )*( vx1 - vx2 ) - cu_sq); + (dist.f[NW ])[knw ] = c1o54* (drho + c3o1*(-vx1 + vx2 ) + c9o2*(-vx1 + vx2 )*(-vx1 + vx2 ) - cu_sq); + (dist.f[TE ])[kte ] = c1o54* (drho + c3o1*( vx1 + vx3) + c9o2*( vx1 + vx3)*( vx1 + vx3) - cu_sq); + (dist.f[BW ])[kbw ] = c1o54* (drho + c3o1*(-vx1 - vx3) + c9o2*(-vx1 - vx3)*(-vx1 - vx3) - cu_sq); + (dist.f[BE ])[kbe ] = c1o54* (drho + c3o1*( vx1 - vx3) + c9o2*( vx1 - vx3)*( vx1 - vx3) - cu_sq); + (dist.f[TW ])[ktw ] = c1o54* (drho + c3o1*(-vx1 + vx3) + c9o2*(-vx1 + vx3)*(-vx1 + vx3) - cu_sq); + (dist.f[TN ])[ktn ] = c1o54* (drho + c3o1*( vx2 + vx3) + c9o2*( vx2 + vx3)*( vx2 + vx3) - cu_sq); + (dist.f[BS ])[kbs ] = c1o54* (drho + c3o1*( - vx2 - vx3) + c9o2*( - vx2 - vx3)*( - vx2 - vx3) - cu_sq); + (dist.f[BN ])[kbn ] = c1o54* (drho + c3o1*( vx2 - vx3) + c9o2*( vx2 - vx3)*( vx2 - vx3) - cu_sq); + (dist.f[TS ])[kts ] = c1o54* (drho + c3o1*( - vx2 + vx3) + c9o2*( - vx2 + vx3)*( - vx2 + vx3) - cu_sq); + (dist.f[TNE ])[ktne ] = c1o216*(drho + c3o1*( vx1 + vx2 + vx3) + c9o2*( vx1 + vx2 + vx3)*( vx1 + vx2 + vx3) - cu_sq); + (dist.f[BSW ])[kbsw ] = c1o216*(drho + c3o1*(-vx1 - vx2 - vx3) + c9o2*(-vx1 - vx2 - vx3)*(-vx1 - vx2 - vx3) - cu_sq); + (dist.f[BNE ])[kbne ] = c1o216*(drho + c3o1*( vx1 + vx2 - vx3) + c9o2*( vx1 + vx2 - vx3)*( vx1 + vx2 - vx3) - cu_sq); + (dist.f[TSW ])[ktsw ] = c1o216*(drho + c3o1*(-vx1 - vx2 + vx3) + c9o2*(-vx1 - vx2 + vx3)*(-vx1 - vx2 + vx3) - cu_sq); + (dist.f[TSE ])[ktse ] = c1o216*(drho + c3o1*( vx1 - vx2 + vx3) + c9o2*( vx1 - vx2 + vx3)*( vx1 - vx2 + vx3) - cu_sq); + (dist.f[BNW ])[kbnw ] = c1o216*(drho + c3o1*(-vx1 + vx2 - vx3) + c9o2*(-vx1 + vx2 - vx3)*(-vx1 + vx2 - vx3) - cu_sq); + (dist.f[BSE ])[kbse ] = c1o216*(drho + c3o1*( vx1 - vx2 - vx3) + c9o2*( vx1 - vx2 - vx3)*( vx1 - vx2 - vx3) - cu_sq); + (dist.f[TNW ])[ktnw ] = c1o216*(drho + c3o1*(-vx1 + vx2 + vx3) + c9o2*(-vx1 + vx2 + vx3)*(-vx1 + vx2 + vx3) - cu_sq); } } diff --git a/src/gpu/VirtualFluids_GPU/GPU/InitAdvectionDiffusion27.cu b/src/gpu/VirtualFluids_GPU/GPU/InitAdvectionDiffusion27.cu index d52d7f8e2..c49c267e2 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/InitAdvectionDiffusion27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/InitAdvectionDiffusion27.cu @@ -76,63 +76,63 @@ extern "C" __global__ void InitAD( Distributions27 distAD; if (isEvenTimestep) { - distAD.f[dirE ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirW ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[E ] = &distributionsAD[E *size_Mat]; + distAD.f[W ] = &distributionsAD[W *size_Mat]; + distAD.f[N ] = &distributionsAD[N *size_Mat]; + distAD.f[S ] = &distributionsAD[S *size_Mat]; + distAD.f[T ] = &distributionsAD[T *size_Mat]; + distAD.f[B ] = &distributionsAD[B *size_Mat]; + distAD.f[NE ] = &distributionsAD[NE *size_Mat]; + distAD.f[SW ] = &distributionsAD[SW *size_Mat]; + distAD.f[SE ] = &distributionsAD[SE *size_Mat]; + distAD.f[NW ] = &distributionsAD[NW *size_Mat]; + distAD.f[TE ] = &distributionsAD[TE *size_Mat]; + distAD.f[BW ] = &distributionsAD[BW *size_Mat]; + distAD.f[BE ] = &distributionsAD[BE *size_Mat]; + distAD.f[TW ] = &distributionsAD[TW *size_Mat]; + distAD.f[TN ] = &distributionsAD[TN *size_Mat]; + distAD.f[BS ] = &distributionsAD[BS *size_Mat]; + distAD.f[BN ] = &distributionsAD[BN *size_Mat]; + distAD.f[TS ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[TNE ] = &distributionsAD[TNE *size_Mat]; + distAD.f[TSW ] = &distributionsAD[TSW *size_Mat]; + distAD.f[TSE ] = &distributionsAD[TSE *size_Mat]; + distAD.f[TNW ] = &distributionsAD[TNW *size_Mat]; + distAD.f[BNE ] = &distributionsAD[BNE *size_Mat]; + distAD.f[BSW ] = &distributionsAD[BSW *size_Mat]; + distAD.f[BSE ] = &distributionsAD[BSE *size_Mat]; + distAD.f[BNW ] = &distributionsAD[BNW *size_Mat]; } else { - distAD.f[dirW ] = &distributionsAD[dirE *size_Mat]; - distAD.f[dirE ] = &distributionsAD[dirW *size_Mat]; - distAD.f[dirS ] = &distributionsAD[dirN *size_Mat]; - distAD.f[dirN ] = &distributionsAD[dirS *size_Mat]; - distAD.f[dirB ] = &distributionsAD[dirT *size_Mat]; - distAD.f[dirT ] = &distributionsAD[dirB *size_Mat]; - distAD.f[dirSW ] = &distributionsAD[dirNE *size_Mat]; - distAD.f[dirNE ] = &distributionsAD[dirSW *size_Mat]; - distAD.f[dirNW ] = &distributionsAD[dirSE *size_Mat]; - distAD.f[dirSE ] = &distributionsAD[dirNW *size_Mat]; - distAD.f[dirBW ] = &distributionsAD[dirTE *size_Mat]; - distAD.f[dirTE ] = &distributionsAD[dirBW *size_Mat]; - distAD.f[dirTW ] = &distributionsAD[dirBE *size_Mat]; - distAD.f[dirBE ] = &distributionsAD[dirTW *size_Mat]; - distAD.f[dirBS ] = &distributionsAD[dirTN *size_Mat]; - distAD.f[dirTN ] = &distributionsAD[dirBS *size_Mat]; - distAD.f[dirTS ] = &distributionsAD[dirBN *size_Mat]; - distAD.f[dirBN ] = &distributionsAD[dirTS *size_Mat]; - distAD.f[dirREST] = &distributionsAD[dirREST*size_Mat]; - distAD.f[dirBSW ] = &distributionsAD[dirTNE *size_Mat]; - distAD.f[dirBNE ] = &distributionsAD[dirTSW *size_Mat]; - distAD.f[dirBNW ] = &distributionsAD[dirTSE *size_Mat]; - distAD.f[dirBSE ] = &distributionsAD[dirTNW *size_Mat]; - distAD.f[dirTSW ] = &distributionsAD[dirBNE *size_Mat]; - distAD.f[dirTNE ] = &distributionsAD[dirBSW *size_Mat]; - distAD.f[dirTNW ] = &distributionsAD[dirBSE *size_Mat]; - distAD.f[dirTSE ] = &distributionsAD[dirBNW *size_Mat]; + distAD.f[W ] = &distributionsAD[E *size_Mat]; + distAD.f[E ] = &distributionsAD[W *size_Mat]; + distAD.f[S ] = &distributionsAD[N *size_Mat]; + distAD.f[N ] = &distributionsAD[S *size_Mat]; + distAD.f[B ] = &distributionsAD[T *size_Mat]; + distAD.f[T ] = &distributionsAD[B *size_Mat]; + distAD.f[SW ] = &distributionsAD[NE *size_Mat]; + distAD.f[NE ] = &distributionsAD[SW *size_Mat]; + distAD.f[NW ] = &distributionsAD[SE *size_Mat]; + distAD.f[SE ] = &distributionsAD[NW *size_Mat]; + distAD.f[BW ] = &distributionsAD[TE *size_Mat]; + distAD.f[TE ] = &distributionsAD[BW *size_Mat]; + distAD.f[TW ] = &distributionsAD[BE *size_Mat]; + distAD.f[BE ] = &distributionsAD[TW *size_Mat]; + distAD.f[BS ] = &distributionsAD[TN *size_Mat]; + distAD.f[TN ] = &distributionsAD[BS *size_Mat]; + distAD.f[TS ] = &distributionsAD[BN *size_Mat]; + distAD.f[BN ] = &distributionsAD[TS *size_Mat]; + distAD.f[REST] = &distributionsAD[REST*size_Mat]; + distAD.f[BSW ] = &distributionsAD[TNE *size_Mat]; + distAD.f[BNE ] = &distributionsAD[TSW *size_Mat]; + distAD.f[BNW ] = &distributionsAD[TSE *size_Mat]; + distAD.f[BSE ] = &distributionsAD[TNW *size_Mat]; + distAD.f[TSW ] = &distributionsAD[BNE *size_Mat]; + distAD.f[TNE ] = &distributionsAD[BSW *size_Mat]; + distAD.f[TNW ] = &distributionsAD[BSE *size_Mat]; + distAD.f[TSE ] = &distributionsAD[BNW *size_Mat]; } ////////////////////////////////////////////////////////////////////////// //! - Set local velocities and concetration @@ -176,33 +176,33 @@ extern "C" __global__ void InitAD( //! real cu_sq = c3o2*(vx1*vx1 + vx2*vx2 + vx3*vx3); - (distAD.f[dirREST])[kzero] = c8o27 * conc * (c1o1 - cu_sq); - (distAD.f[dirE ])[ke ] = c2o27 * conc * (c1o1 + c3o1 * ( vx1 ) + c9o2 * ( vx1 ) * ( vx1 ) - cu_sq); - (distAD.f[dirW ])[kw ] = c2o27 * conc * (c1o1 + c3o1 * (-vx1 ) + c9o2 * (-vx1 ) * (-vx1 ) - cu_sq); - (distAD.f[dirN ])[kn ] = c2o27 * conc * (c1o1 + c3o1 * ( vx2 ) + c9o2 * ( vx2 ) * ( vx2 ) - cu_sq); - (distAD.f[dirS ])[ks ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx2 ) + c9o2 * ( - vx2 ) * ( - vx2 ) - cu_sq); - (distAD.f[dirT ])[kt ] = c2o27 * conc * (c1o1 + c3o1 * ( vx3) + c9o2 * ( vx3) * ( vx3) - cu_sq); - (distAD.f[dirB ])[kb ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx3) + c9o2 * ( - vx3) * ( - vx3) - cu_sq); - (distAD.f[dirNE ])[kne ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx2 ) + c9o2 * ( vx1 + vx2 ) * ( vx1 + vx2 ) - cu_sq); - (distAD.f[dirSW ])[ksw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx2 ) + c9o2 * (-vx1 - vx2 ) * (-vx1 - vx2 ) - cu_sq); - (distAD.f[dirSE ])[kse ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx2 ) + c9o2 * ( vx1 - vx2 ) * ( vx1 - vx2 ) - cu_sq); - (distAD.f[dirNW ])[knw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx2 ) + c9o2 * (-vx1 + vx2 ) * (-vx1 + vx2 ) - cu_sq); - (distAD.f[dirTE ])[kte ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx3) + c9o2 * ( vx1 + vx3) * ( vx1 + vx3) - cu_sq); - (distAD.f[dirBW ])[kbw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx3) + c9o2 * (-vx1 - vx3) * (-vx1 - vx3) - cu_sq); - (distAD.f[dirBE ])[kbe ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx3) + c9o2 * ( vx1 - vx3) * ( vx1 - vx3) - cu_sq); - (distAD.f[dirTW ])[ktw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx3) + c9o2 * (-vx1 + vx3) * (-vx1 + vx3) - cu_sq); - (distAD.f[dirTN ])[ktn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 + vx3) + c9o2 * ( vx2 + vx3) * ( vx2 + vx3) - cu_sq); - (distAD.f[dirBS ])[kbs ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 - vx3) + c9o2 * ( - vx2 - vx3) * ( - vx2 - vx3) - cu_sq); - (distAD.f[dirBN ])[kbn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 - vx3) + c9o2 * ( vx2 - vx3) * ( vx2 - vx3) - cu_sq); - (distAD.f[dirTS ])[kts ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 + vx3) + c9o2 * ( - vx2 + vx3) * ( - vx2 + vx3) - cu_sq); - (distAD.f[dirTNE ])[ktne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 + vx3) + c9o2 * ( vx1 + vx2 + vx3) * ( vx1 + vx2 + vx3) - cu_sq); - (distAD.f[dirBSW ])[kbsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 - vx3) + c9o2 * (-vx1 - vx2 - vx3) * (-vx1 - vx2 - vx3) - cu_sq); - (distAD.f[dirBNE ])[kbne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 - vx3) + c9o2 * ( vx1 + vx2 - vx3) * ( vx1 + vx2 - vx3) - cu_sq); - (distAD.f[dirTSW ])[ktsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 + vx3) + c9o2 * (-vx1 - vx2 + vx3) * (-vx1 - vx2 + vx3) - cu_sq); - (distAD.f[dirTSE ])[ktse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 + vx3) + c9o2 * ( vx1 - vx2 + vx3) * ( vx1 - vx2 + vx3) - cu_sq); - (distAD.f[dirBNW ])[kbnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 - vx3) + c9o2 * (-vx1 + vx2 - vx3) * (-vx1 + vx2 - vx3) - cu_sq); - (distAD.f[dirBSE ])[kbse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 - vx3) + c9o2 * ( vx1 - vx2 - vx3) * ( vx1 - vx2 - vx3) - cu_sq); - (distAD.f[dirTNW ])[ktnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 + vx3) + c9o2 * (-vx1 + vx2 + vx3) * (-vx1 + vx2 + vx3) - cu_sq); + (distAD.f[REST])[kzero] = c8o27 * conc * (c1o1 - cu_sq); + (distAD.f[E ])[ke ] = c2o27 * conc * (c1o1 + c3o1 * ( vx1 ) + c9o2 * ( vx1 ) * ( vx1 ) - cu_sq); + (distAD.f[W ])[kw ] = c2o27 * conc * (c1o1 + c3o1 * (-vx1 ) + c9o2 * (-vx1 ) * (-vx1 ) - cu_sq); + (distAD.f[N ])[kn ] = c2o27 * conc * (c1o1 + c3o1 * ( vx2 ) + c9o2 * ( vx2 ) * ( vx2 ) - cu_sq); + (distAD.f[S ])[ks ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx2 ) + c9o2 * ( - vx2 ) * ( - vx2 ) - cu_sq); + (distAD.f[T ])[kt ] = c2o27 * conc * (c1o1 + c3o1 * ( vx3) + c9o2 * ( vx3) * ( vx3) - cu_sq); + (distAD.f[B ])[kb ] = c2o27 * conc * (c1o1 + c3o1 * ( - vx3) + c9o2 * ( - vx3) * ( - vx3) - cu_sq); + (distAD.f[NE ])[kne ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx2 ) + c9o2 * ( vx1 + vx2 ) * ( vx1 + vx2 ) - cu_sq); + (distAD.f[SW ])[ksw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx2 ) + c9o2 * (-vx1 - vx2 ) * (-vx1 - vx2 ) - cu_sq); + (distAD.f[SE ])[kse ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx2 ) + c9o2 * ( vx1 - vx2 ) * ( vx1 - vx2 ) - cu_sq); + (distAD.f[NW ])[knw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx2 ) + c9o2 * (-vx1 + vx2 ) * (-vx1 + vx2 ) - cu_sq); + (distAD.f[TE ])[kte ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 + vx3) + c9o2 * ( vx1 + vx3) * ( vx1 + vx3) - cu_sq); + (distAD.f[BW ])[kbw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 - vx3) + c9o2 * (-vx1 - vx3) * (-vx1 - vx3) - cu_sq); + (distAD.f[BE ])[kbe ] = c1o54 * conc * (c1o1 + c3o1 * ( vx1 - vx3) + c9o2 * ( vx1 - vx3) * ( vx1 - vx3) - cu_sq); + (distAD.f[TW ])[ktw ] = c1o54 * conc * (c1o1 + c3o1 * (-vx1 + vx3) + c9o2 * (-vx1 + vx3) * (-vx1 + vx3) - cu_sq); + (distAD.f[TN ])[ktn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 + vx3) + c9o2 * ( vx2 + vx3) * ( vx2 + vx3) - cu_sq); + (distAD.f[BS ])[kbs ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 - vx3) + c9o2 * ( - vx2 - vx3) * ( - vx2 - vx3) - cu_sq); + (distAD.f[BN ])[kbn ] = c1o54 * conc * (c1o1 + c3o1 * ( vx2 - vx3) + c9o2 * ( vx2 - vx3) * ( vx2 - vx3) - cu_sq); + (distAD.f[TS ])[kts ] = c1o54 * conc * (c1o1 + c3o1 * ( - vx2 + vx3) + c9o2 * ( - vx2 + vx3) * ( - vx2 + vx3) - cu_sq); + (distAD.f[TNE ])[ktne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 + vx3) + c9o2 * ( vx1 + vx2 + vx3) * ( vx1 + vx2 + vx3) - cu_sq); + (distAD.f[BSW ])[kbsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 - vx3) + c9o2 * (-vx1 - vx2 - vx3) * (-vx1 - vx2 - vx3) - cu_sq); + (distAD.f[BNE ])[kbne ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 + vx2 - vx3) + c9o2 * ( vx1 + vx2 - vx3) * ( vx1 + vx2 - vx3) - cu_sq); + (distAD.f[TSW ])[ktsw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 - vx2 + vx3) + c9o2 * (-vx1 - vx2 + vx3) * (-vx1 - vx2 + vx3) - cu_sq); + (distAD.f[TSE ])[ktse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 + vx3) + c9o2 * ( vx1 - vx2 + vx3) * ( vx1 - vx2 + vx3) - cu_sq); + (distAD.f[BNW ])[kbnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 - vx3) + c9o2 * (-vx1 + vx2 - vx3) * (-vx1 + vx2 - vx3) - cu_sq); + (distAD.f[BSE ])[kbse ] = c1o216 * conc * (c1o1 + c3o1 * ( vx1 - vx2 - vx3) + c9o2 * ( vx1 - vx2 - vx3) * ( vx1 - vx2 - vx3) - cu_sq); + (distAD.f[TNW ])[ktnw ] = c1o216 * conc * (c1o1 + c3o1 * (-vx1 + vx2 + vx3) + c9o2 * (-vx1 + vx2 + vx3) * (-vx1 + vx2 + vx3) - cu_sq); } } diff --git a/src/gpu/VirtualFluids_GPU/LBM/D3Q27.h b/src/gpu/VirtualFluids_GPU/LBM/D3Q27.h index ab5d960b5..80c6cfa3b 100644 --- a/src/gpu/VirtualFluids_GPU/LBM/D3Q27.h +++ b/src/gpu/VirtualFluids_GPU/LBM/D3Q27.h @@ -33,35 +33,35 @@ #ifndef _LB_D3Q27_H_ #define _LB_D3Q27_H_ -static constexpr int dirSTART = 0; -static constexpr int dirEND = 26; +static constexpr int STARTDIR = 0; +static constexpr int ENDDIR = 26; -static constexpr int dirE = 0; -static constexpr int dirW = 1; -static constexpr int dirN = 2; -static constexpr int dirS = 3; -static constexpr int dirT = 4; -static constexpr int dirB = 5; -static constexpr int dirNE = 6; -static constexpr int dirSW = 7; -static constexpr int dirSE = 8; -static constexpr int dirNW = 9; -static constexpr int dirTE = 10; -static constexpr int dirBW = 11; -static constexpr int dirBE = 12; -static constexpr int dirTW = 13; -static constexpr int dirTN = 14; -static constexpr int dirBS = 15; -static constexpr int dirBN = 16; -static constexpr int dirTS = 17; -static constexpr int dirTNE = 18; -static constexpr int dirTNW = 19; -static constexpr int dirTSE = 20; -static constexpr int dirTSW = 21; -static constexpr int dirBNE = 22; -static constexpr int dirBNW = 23; -static constexpr int dirBSE = 24; -static constexpr int dirBSW = 25; +static constexpr int E = 0; +static constexpr int W = 1; +static constexpr int N = 2; +static constexpr int S = 3; +static constexpr int T = 4; +static constexpr int B = 5; +static constexpr int NE = 6; +static constexpr int SW = 7; +static constexpr int SE = 8; +static constexpr int NW = 9; +static constexpr int TE = 10; +static constexpr int BW = 11; +static constexpr int BE = 12; +static constexpr int TW = 13; +static constexpr int TN = 14; +static constexpr int BS = 15; +static constexpr int BN = 16; +static constexpr int TS = 17; +static constexpr int TNE = 18; +static constexpr int TNW = 19; +static constexpr int TSE = 20; +static constexpr int TSW = 21; +static constexpr int BNE = 22; +static constexpr int BNW = 23; +static constexpr int BSE = 24; +static constexpr int BSW = 25; static constexpr int dirREST = 26; -- GitLab