diff --git a/src/gpu/VirtualFluids_GPU/AdvectionDiffusion/AdvectionDiffusion.cpp b/src/gpu/VirtualFluids_GPU/AdvectionDiffusion/AdvectionDiffusion.cpp
index f466cb5c0a21fdf7119cf7bc682d2b432a5e3a96..27278684a1875c41b855076d4ee702d721d9a9fe 100644
--- a/src/gpu/VirtualFluids_GPU/AdvectionDiffusion/AdvectionDiffusion.cpp
+++ b/src/gpu/VirtualFluids_GPU/AdvectionDiffusion/AdvectionDiffusion.cpp
@@ -40,11 +40,11 @@ void initAD(SPtr<Parameter> para)
{
//////////////////////////////////////////////////////////////////////////
// calculation of omega for diffusivity
- para->getParD()->omegaD = (real)2.0 / ((real)6.0 * para->getParD()->diffusivity + (real)1.0);
+ para->getParD()->omegaDiffusivity = (real)2.0 / ((real)6.0 * para->getParD()->diffusivity + (real)1.0);
//////////////////////////////////////////////////////////////////////////
para->getParD()->isEvenTimestep = true;
//////////////////////////////////////////////////////////////////////////
- InitADdevice(
+ InitADDev27(
para->getParD()->numberofthreads,
para->getParD()->neighborX,
para->getParD()->neighborY,
@@ -60,7 +60,7 @@ void initAD(SPtr<Parameter> para)
//////////////////////////////////////////////////////////////////////////
para->getParD()->isEvenTimestep = false;
//////////////////////////////////////////////////////////////////////////
- InitADdevice(
+ InitADDev27(
para->getParD()->numberofthreads,
para->getParD()->neighborX,
para->getParD()->neighborY,
@@ -110,11 +110,11 @@ void setInitialNodeValuesAD(SPtr<Parameter> para, SPtr<CudaMemoryManager> cudaMe
}
////////////////////////////////////////////////////////////////////////////////
-void calcAD(SPtr<Parameter> para)
+void runADcollisionKernel(SPtr<Parameter> para)
{
FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
para->getParD()->numberofthreads,
- para->getParD()->omegaD,
+ para->getParD()->omegaDiffusivity,
para->getParD()->typeOfGridNode,
para->getParD()->neighborX,
para->getParD()->neighborY,
@@ -124,19 +124,21 @@ void calcAD(SPtr<Parameter> para)
para->getParD()->numberOfNodes,
para->getParD()->forcing,
para->getParD()->isEvenTimestep);
+}
+void runADslipBCKernel(){
if (para->getParD()->numberOfSlipBCnodes > 1) {
ADSlipVelDevComp(
para->getParD()->numberofthreads,
- para->getParD()->slipBC.NormalX,
- para->getParD()->slipBC.NormalY,
- para->getParD()->slipBC.NormalZ,
+ para->getParD()->slipBC.normalX,
+ para->getParD()->slipBC.normalY,
+ para->getParD()->slipBC.normalZ,
para->getParD()->distributions.f[0],
para->getParD()->distributionsAD.f[0],
para->getParD()->slipBC.k,
para->getParD()->slipBC.q27[0],
para->getParD()->numberOfSlipBCnodes,
- para->getParD()->omegaD,
+ para->getParD()->omegaDiffusivity,
para->getParD()->neighborX,
para->getParD()->neighborY,
para->getParD()->neighborZ,
@@ -145,6 +147,7 @@ void calcAD(SPtr<Parameter> para)
}
}
+
////////////////////////////////////////////////////////////////////////////////
void printAD(SPtr<Parameter> para, SPtr<CudaMemoryManager> cudaMemoryManager)
{
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index a68e8543f8673329f51be9de97d108e6eba4233f..7a89a5aa8edfba72c5b6298d26a9ab9825c259d3 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -105,9 +105,9 @@ void GridGenerator::allocArrays_BoundaryValues()
cudaMemoryManager->cudaAllocSlipBC();
builder->getSlipValues(
- para->getParH()->slipBC.NormalX,
- para->getParH()->slipBC.NormalY,
- para->getParH()->slipBC.NormalY,
+ para->getParH()->slipBC.normalX,
+ para->getParH()->slipBC.normalY,
+ para->getParH()->slipBC.normalY,
para->getParH()->slipBC.k, 0);
cudaMemoryManager->cudaCopySlipBC();
diff --git a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu
index ea4d643b397b6fc04bf973383cf4d5de341bdb89..8b99bd07f9928aab7bb8b62b6757775b59ce2537 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/AdvectionDiffusion27chim.cu
@@ -67,7 +67,7 @@ inline __device__ void backwardChimera(real &mfa, real &mfb, real &mfc, real vv,
////////////////////////////////////////////////////////////////////////////////
extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device_Kernel(
- real omegaD,
+ real omegaDiffusivity,
uint* typeOfGridNode,
uint* neighborX,
uint* neighborY,
@@ -347,7 +347,7 @@ extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device
real vy2 = vvy*vvy;
real vz2 = vvz*vvz;
////////////////////////////////////////////////////////////////////////////////////
- //real omegaD = c2o1 / (c6o1 * diffusivity + c1o1);
+ //real omegaDiffusivity = c2o1 / (c6o1 * diffusivity + c1o1);
////////////////////////////////////////////////////////////////////////////////////
//! - Chimera transform from distributions to central moments as defined in Eq. (43)-(45) in \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>
@@ -406,9 +406,9 @@ extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device
real Mccb = mfccb - mfaab*c1o9;
// collision of 1st order moments
- mfbaa *= c1o1 - omegaD;
- mfaba *= c1o1 - omegaD;
- mfaab *= c1o1 - omegaD;
+ mfbaa *= c1o1 - omegaDiffusivity;
+ mfaba *= c1o1 - omegaDiffusivity;
+ mfaab *= c1o1 - omegaDiffusivity;
// equilibration of 3rd order moments
Mabc = c0o1;
diff --git a/src/gpu/VirtualFluids_GPU/GPU/BCsAdvectionDiffusion27.cu b/src/gpu/VirtualFluids_GPU/GPU/BCsAdvectionDiffusion27.cu
index b1fb3456117f7dd65de5439fd6b987a7d984c92e..ad7fcf548d46c425f80aa16f87d4be3ab0458d27 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/BCsAdvectionDiffusion27.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/BCsAdvectionDiffusion27.cu
@@ -94,12 +94,12 @@ inline __device__ void backwardChimera(real &mfa, real &mfb, real &mfc, real vv,
}
////////////////////////////////////////////////////////////////////////////////
-inline __device__ real calcDistributionBC_AD_interpol(real q, real weight, real v, real v_sq, real f, real finf, real omegaD, real jTangential, real concentration) {
+inline __device__ real calcDistributionBC_AD_interpol(real q, real weight, real v, real v_sq, real f, real finf, real omegaDiffusivity, real jTangential, real concentration) {
real feq = weight * concentration * (c1o1 + c3o1 * v + c9o2 * v * v * concentration - v_sq * concentration);
- return (c1o1 - q) / (c1o1 + q) * ((f - feq * omegaD) / (c1o1 - omegaD)) + (q * (f + finf) - c6o1 * weight * (jTangential)) / (c1o1 + q);
+ return (c1o1 - q) / (c1o1 + q) * ((f - feq * omegaDiffusivity) / (c1o1 - omegaDiffusivity)) + (q * (f + finf) - c6o1 * weight * (jTangential)) / (c1o1 + q);
}
////////////////////////////////////////////////////////////////////////////////
-inline __device__ real calcDistributionBC_AD(real q, real weight, real v, real v_sq, real f, real finf, real omegaD, real jTangential, real concentration) {
+inline __device__ real calcDistributionBC_AD(real q, real weight, real v, real v_sq, real f, real finf, real omegaDiffusivity, real jTangential, real concentration) {
return f - c6o1 * weight * jTangential;
}
@@ -107,15 +107,15 @@ inline __device__ real calcDistributionBC_AD(real q, real weight, real v, real v
// has to be excecuted before Fluid BCs
//////////////////////////////////////////////////////////////////////////////
extern "C" __global__ void AD_SlipVelDeviceComp(
- real *NormalX,
- real *NormalY,
- real *NormalZ,
+ real *normalX,
+ real *normalY,
+ real *normalZ,
real *distributions,
real *distributionsAD,
int *QindexArray,
real *Qarrays,
uint numberOfQs,
- real omegaD,
+ real omegaDiffusivity,
uint* neighborX,
uint* neighborY,
uint* neighborZ,
@@ -259,9 +259,9 @@ extern "C" __global__ void AD_SlipVelDeviceComp(
if (k < numberOfQs)
{
////////////////////////////////////////////////////////////////////////////////
- real NormX = NormalX[k];
- real NormY = NormalY[k];
- real NormZ = NormalZ[k];
+ real NormX = normalX[k];
+ real NormY = normalY[k];
+ real NormZ = normalZ[k];
////////////////////////////////////////////////////////////////////////////////
real* q_dirE, * q_dirW, * q_dirN, * q_dirS, * q_dirT, * q_dirB,
* q_dirNE, * q_dirSW, * q_dirSE, * q_dirNW, * q_dirTE, * q_dirBW,
@@ -484,9 +484,9 @@ extern "C" __global__ void AD_SlipVelDeviceComp(
(-(f_BN - f_TS) + (f_TN - f_BS)) + ((f_TE - f_BW) - (f_BE - f_TW)) +
(f_T - f_B)) - (vx3 * concentration);
- //jx1 *= (c2o1 - omegaD) / (c2o1 - c2o1 * omegaD);
- //jx2 *= (c2o1 - omegaD) / (c2o1 - c2o1 * omegaD);
- //jx3 *= (c2o1 - omegaD) / (c2o1 - c2o1 * omegaD);
+ //jx1 *= (c2o1 - omegaDiffusivity) / (c2o1 - c2o1 * omegaDiffusivity);
+ //jx2 *= (c2o1 - omegaDiffusivity) / (c2o1 - c2o1 * omegaDiffusivity);
+ //jx3 *= (c2o1 - omegaDiffusivity) / (c2o1 - c2o1 * omegaDiffusivity);
real NormJ = jx1 * NormX + jx2 * NormY + jx3 * NormZ;
@@ -495,32 +495,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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -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, omegaD, 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, omegaD, -jTan1+jTan2+jTan3, concentration); }
+ 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); }
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
}
}
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu b/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu
index 4067a702a6fdf76e073c1a8c26833d2e338bf227..3b4db0c2e87a986a9442cf87d62106ec62da9c7f 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/CalcConc27.cu
@@ -54,9 +54,9 @@ extern "C" __global__ void CalcConc27(
////////////////////////////////////////////////////////////////////////////////
//! - Get node index coordinates from thredIdx, blockIdx, blockDim and gridDim.
//!
- const unsigned x = threadIdx.x; // Globaler x-Index
- const unsigned y = blockIdx.x; // Globaler y-Index
- const unsigned z = blockIdx.y; // Globaler z-Index
+ const unsigned x = threadIdx.x; // global x-index
+ const unsigned y = blockIdx.x; // global y-index
+ const unsigned z = blockIdx.y; // global z-index
const unsigned nx = blockDim.x;
const unsigned ny = gridDim.x;
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 364843deb94fd6730940055f5cba10179f057265..9766124ba2b22440ac7527d2e54a488ed52bf194 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -149,16 +149,16 @@ void CudaMemoryManager::cudaAllocSlipBC()
//Host
checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.q27[0]), parameter->getD3Qxx() * mem_size_slip_BC_REAL));
checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.k), mem_size_slip_BC_INT));
- checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.NormalX), mem_size_slip_BC_REAL));
- checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.NormalY), mem_size_slip_BC_REAL));
- checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.NormalZ), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.normalX), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.normalY), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMallocHost((void**)&(parameter->getParH()->slipBC.normalZ), mem_size_slip_BC_REAL));
//Device
checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.q27[0]), parameter->getD3Qxx() * mem_size_slip_BC_REAL));
checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.k), mem_size_slip_BC_INT));
- checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.NormalX), mem_size_slip_BC_REAL));
- checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.NormalY), mem_size_slip_BC_REAL));
- checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.NormalZ), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.normalX), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.normalY), mem_size_slip_BC_REAL));
+ checkCudaErrors(cudaMalloc((void**)&(parameter->getParD()->slipBC.normalZ), mem_size_slip_BC_REAL));
//////////////////////////////////////////////////////////////////////////
double tmp = (double)mem_size_slip_BC_INT + 4. * (double)mem_size_slip_BC_REAL + (double)parameter->getD3Qxx() * (double)mem_size_slip_BC_REAL;
@@ -171,18 +171,18 @@ void CudaMemoryManager::cudaCopySlipBC()
checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.q27[0], parameter->getParH()->slipBC.q27[0], parameter->getD3Qxx() * mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.k, parameter->getParH()->slipBC.k, mem_size_slip_BC_INT, cudaMemcpyHostToDevice));
- checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.NormalX, parameter->getParH()->slipBC.NormalX, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.NormalY, parameter->getParH()->slipBC.NormalY, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.NormalZ, parameter->getParH()->slipBC.NormalZ, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.normalX, parameter->getParH()->slipBC.normalX, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.normalY, parameter->getParH()->slipBC.normalY, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors(cudaMemcpy(parameter->getParD()->slipBC.normalZ, parameter->getParH()->slipBC.normalZ, mem_size_slip_BC_REAL, cudaMemcpyHostToDevice));
}
void CudaMemoryManager::cudaFreeSlipBC()
{
checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.q27[0]));
checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.k));
- checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.NormalX));
- checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.NormalY));
- checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.NormalZ));
+ checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.normalX));
+ checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.normalY));
+ checkCudaErrors(cudaFreeHost(parameter->getParH()->slipBC.normalZ));
}
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
index ac4d43ac785e7ae8134d367a3c30b806e8e556c5..e270871b0d3dc2dd49beb7e944ddd3926391fbe5 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -141,7 +141,7 @@ extern "C" void QVelDevicePlainBB27(
//! \brief Advection Diffusion kernel
extern "C" void FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
uint numberOfThreads,
- real omegaD,
+ real omegaDiffusivity,
uint* typeOfGridNode,
uint* neighborX,
uint* neighborY,
@@ -154,7 +154,7 @@ extern "C" void FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
//////////////////////////////////////////////////////////////////////////
//! \brief initialize the Advection Diffusion distribution functions
-extern "C" void InitADdevice(
+extern "C" void InitADDev27(
uint numberOfThreads,
uint* neighborX,
uint* neighborY,
@@ -185,15 +185,15 @@ extern "C" void CalcConcentration27(
//! \brief defines the behavior of a slip-AD boundary condition
extern "C" void ADSlipVelDevComp(
uint numberOfThreads,
- real * NormalX,
- real * NormalY,
- real * NormalZ,
+ real * normalX,
+ real * normalY,
+ real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfQs,
- real omegaD,
+ real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
index 32cca83ed57dcad805202e0f9e5c0649b5c6e531..dc77c8f12338b01dd216c738aac5e75add351b85 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
@@ -136,7 +136,7 @@ extern "C" __global__ void QVelDevPlainBB27(
//////////////////////////////////////////////////////////////////////////
//! \brief \ref Advection_Diffusion_Device_Kernel : Factorized central moments for Advection Diffusion Equation
extern "C" __global__ void Factorized_Central_Moments_Advection_Diffusion_Device_Kernel(
- real omegaD,
+ real omegaDiffusivity,
uint* typeOfGridNode,
uint* neighborX,
uint* neighborY,
@@ -177,15 +177,15 @@ extern "C" __global__ void CalcConc27(
//////////////////////////////////////////////////////////////////////////
//! \brief \ref AD_SlipVelDeviceComp : device function for the slip-AD boundary condition
extern "C" __global__ void AD_SlipVelDeviceComp(
- real * NormalX,
- real * NormalY,
- real * NormalZ,
+ real * normalX,
+ real * normalY,
+ real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfQs,
- real omegaD,
+ real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index 784793b54018f322fc48cef5660a618fad1ecd99..34dac34640ff926549ecdb5cf0b831155fb057ce 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -254,7 +254,7 @@ extern "C" void QVelDevicePlainBB27(
//////////////////////////////////////////////////////////////////////////
extern "C" void FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
uint numberOfThreads,
- real omegaD,
+ real omegaDiffusivity,
uint* typeOfGridNode,
uint* neighborX,
uint* neighborY,
@@ -270,7 +270,7 @@ extern "C" void FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
dim3 threads(numberOfThreads, 1, 1);
Factorized_Central_Moments_Advection_Diffusion_Device_Kernel <<< grid, threads >>> (
- omegaD,
+ omegaDiffusivity,
typeOfGridNode,
neighborX,
neighborY,
@@ -283,7 +283,7 @@ extern "C" void FactorizedCentralMomentsAdvectionDiffusionDeviceKernel(
getLastCudaError("Factorized_Central_Moments_Advection_Diffusion_Device_Kernel execution failed");
}
//////////////////////////////////////////////////////////////////////////
-extern "C" void InitADdevice(
+extern "C" void InitADDev27(
uint numberOfThreads,
uint* neighborX,
uint* neighborY,
@@ -345,15 +345,15 @@ extern "C" void CalcConcentration27(
//////////////////////////////////////////////////////////////////////////
extern "C" void ADSlipVelDevComp(
uint numberOfThreads,
- real * NormalX,
- real * NormalY,
- real * NormalZ,
+ real * normalX,
+ real * normalY,
+ real * normalZ,
real * distributions,
real * distributionsAD,
int* QindexArray,
real * Qarrays,
uint numberOfQs,
- real omegaD,
+ real omegaDiffusivity,
uint * neighborX,
uint * neighborY,
uint * neighborZ,
@@ -365,15 +365,15 @@ extern "C" void ADSlipVelDevComp(
dim3 threads(numberOfThreads, 1, 1);
AD_SlipVelDeviceComp << < gridQ, threads >> > (
- NormalX,
- NormalY,
- NormalZ,
+ normalX,
+ normalY,
+ normalZ,
distributions,
distributionsAD,
QindexArray,
Qarrays,
numberOfQs,
- omegaD,
+ omegaDiffusivity,
neighborX,
neighborY,
neighborZ,
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index 133ad8d672194b0c44382612d46e85c518c6467c..cc9aafcb3e40b62b39f997ba526e5f9dd1eb53a7 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -55,7 +55,7 @@ typedef struct QforBC{
real* q27[27];
int kArray;
real *Vx, *Vy, *Vz;
- real *NormalX, *NormalY, *NormalZ;
+ real *normalX, *normalY, *normalZ;
} QforBoundaryConditions;
#endif
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 3e212482de466c4c4451d9d02d2ce6433a02edce..6018486ef06cc59f0f48c55ad8137684687be229 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -89,7 +89,7 @@ struct ParameterStruct{
//! \brief number of lattice nodes for the velocity boundary condition
uint numberOfgeometryBCnodes;
//////////////////////////////////////////////////////////////////////////
- //! \brief sets the forcing uniform on every fluid node in all three space dimensions
+ //! \brief sets the forcing uniform on every fluid node in all three space dimensions
real *forcing;
//////////////////////////////////////////////////////////////////////////
@@ -98,7 +98,7 @@ struct ParameterStruct{
//! \brief stores the diffusivity
real diffusivity;
//! \brief stores the value for omega (for the diffusivity)
- real omegaD;
+ real omegaDiffusivity;
//! \brief stores a field of concentration values
real *concentration;
//! \brief store all distribution functions for the D3Q27 Advection Diffusion field