diff --git a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h b/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h index f7cd7cc1a0dd1fc941547b5d102719a82eef9ca2..c6877cbfeffe5b32c0c2d336e46b02d68cd946a3 100644 --- a/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h +++ b/src/gpu/VirtualFluids_GPU/Factories/BoundaryConditionFactory.h @@ -48,7 +48,7 @@ class Parameter; using boundaryCondition = std::function<void(LBMSimulationParameter *, QforBoundaryConditions *)>; using boundaryConditionWithParameter = std::function<void(Parameter *, QforBoundaryConditions *, const int level)>; -using precursorBoundaryConditionFunc = std::function<void(LBMSimulationParameter *, QforPrecursorBoundaryConditions *, real tRatio, real velocityRatio)>; +using precursorBoundaryConditionFunc = std::function<void(LBMSimulationParameter *, QforPrecursorBoundaryConditions *, real timeRatio, real velocityRatio)>; class BoundaryConditionFactory { diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh index b35e01eb997723eb12f5645857bc230536fe97fe..9ef5057bbf12af887d438e49b974402136fc60c1 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh +++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh @@ -1266,7 +1266,7 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices, real velocityX, real velocityY, real velocityZ, - real tRatio, + real timeRatio, real velocityRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep); @@ -1292,7 +1292,7 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices, real velocityX, real velocityY, real velocityZ, - real tRatio, + real timeRatio, real velocityRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep); @@ -1314,7 +1314,7 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices, real* weightsSB, real* fsLast, real* fsNext, - real tRatio, + real timeRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep); __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices, @@ -1336,7 +1336,7 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices, real* weightsSB, real* fsLast, real* fsNext, - real tRatio, + real timeRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep); diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu index ae51ce98e449bd3251f6b4df2a3d430815892a8a..06628edb5e3d47f5b0ed44ce61ec6da134c551da 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu @@ -3208,7 +3208,7 @@ void VelSchlaffer27( unsigned int numberOfThreads, getLastCudaError("VelSchlaff27 execution failed"); } ////////////////////////////////////////////////////////////////////////// -void QPrecursorDevCompZeroPress(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio) +void QPrecursorDevCompZeroPress(LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real timeRatio, real velocityRatio) { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes); @@ -3220,12 +3220,12 @@ void QPrecursorDevCompZeroPress(LBMSimulationParameter* parameterDevice, QforPre boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB, boundaryCondition->last, boundaryCondition->current, boundaryCondition->velocityX, boundaryCondition->velocityY, boundaryCondition->velocityZ, - tRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); + timeRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); getLastCudaError("QPrecursorDeviceCompZeroPress execution failed"); } ////////////////////////////////////////////////////////////////////////// -void PrecursorDevEQ27( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio) +void PrecursorDevEQ27( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real timeRatio, real velocityRatio) { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes); @@ -3236,12 +3236,12 @@ void PrecursorDevEQ27( LBMSimulationParameter* parameterDevice, QforPrecursorBou boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB, boundaryCondition->last, boundaryCondition->current, boundaryCondition->velocityX, boundaryCondition->velocityY, boundaryCondition->velocityZ, - tRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); + timeRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); getLastCudaError("PrecursorDeviceEQ27 execution failed"); } ////////////////////////////////////////////////////////////////////////// -void PrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio) +void PrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real timeRatio, real velocityRatio) { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes); @@ -3251,13 +3251,13 @@ void PrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPre boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB, boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB, boundaryCondition->last, boundaryCondition->current, - tRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); + timeRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); getLastCudaError("QPrecursorDeviceCompZeroPress execution failed"); } ////////////////////////////////////////////////////////////////////////// -void QPrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real tRatio, real velocityRatio) +void QPrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPrecursorBoundaryConditions* boundaryCondition, real timeRatio, real velocityRatio) { vf::cuda::CudaGrid grid = vf::cuda::CudaGrid(parameterDevice->numberofthreads, boundaryCondition->numberOfBCnodes); @@ -3267,7 +3267,7 @@ void QPrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPr boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB, boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB, boundaryCondition->last, boundaryCondition->current, - tRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); + timeRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep); getLastCudaError("QPrecursorDeviceCompZeroPress execution failed"); } diff --git a/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu b/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu index 9c7958cdc6bc783ba8aadcd0e19877f491fad085..f089a32955af615bfa5744857e1c7d76e2bb42d3 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu @@ -32,7 +32,7 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices, real velocityX, real velocityY, real velocityZ, - real tRatio, + real timeRatio, real velocityRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep) @@ -89,9 +89,9 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices, } // if(k==16300)s printf("%f %f %f\n", vxLastInterpd, vyLastInterpd, vzLastInterpd); - real VeloX = (velocityX + (1.f-tRatio)*vxLastInterpd + tRatio*vxNextInterpd)/velocityRatio; - real VeloY = (velocityY + (1.f-tRatio)*vyLastInterpd + tRatio*vyNextInterpd)/velocityRatio; - real VeloZ = (velocityZ + (1.f-tRatio)*vzLastInterpd + tRatio*vzNextInterpd)/velocityRatio; + real VeloX = (velocityX + (1.f-timeRatio)*vxLastInterpd + timeRatio*vxNextInterpd)/velocityRatio; + real VeloY = (velocityY + (1.f-timeRatio)*vyLastInterpd + timeRatio*vyNextInterpd)/velocityRatio; + real VeloZ = (velocityZ + (1.f-timeRatio)*vzLastInterpd + timeRatio*vzNextInterpd)/velocityRatio; // From here on just a copy of QVelDeviceCompZeroPress //////////////////////////////////////////////////////////////////////////////// @@ -445,7 +445,7 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices, real velocityX, real velocityY, real velocityZ, - real tRatio, + real timeRatio, real velocityRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep) @@ -502,9 +502,9 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices, } // if(k==16300) printf("%f %f %f\n", vxLastInterpd, vyLastInterpd, vzLastInterpd); - real VeloX = (velocityX + (1.f-tRatio)*vxLastInterpd + tRatio*vxNextInterpd)/velocityRatio; - real VeloY = (velocityY + (1.f-tRatio)*vyLastInterpd + tRatio*vyNextInterpd)/velocityRatio; - real VeloZ = (velocityZ + (1.f-tRatio)*vzLastInterpd + tRatio*vzNextInterpd)/velocityRatio; + real VeloX = (velocityX + (1.f-timeRatio)*vxLastInterpd + timeRatio*vxNextInterpd)/velocityRatio; + real VeloY = (velocityY + (1.f-timeRatio)*vyLastInterpd + timeRatio*vyNextInterpd)/velocityRatio; + real VeloZ = (velocityZ + (1.f-timeRatio)*vzLastInterpd + timeRatio*vzNextInterpd)/velocityRatio; // From here on just a copy of QVelDeviceCompZeroPress //////////////////////////////////////////////////////////////////////////////// @@ -666,7 +666,7 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices, real* weightsSB, real* fsLast, real* fsNext, - real tRatio, + real timeRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep) { @@ -793,15 +793,15 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices, unsigned int ktne = KQK; // unsigned int kbsw = neighborZ[ksw]; - dist.f[DIR_P00][ke] = f0LastInterp*(1.f-tRatio) + f0NextInterp*tRatio; - dist.f[DIR_PP0][kne] = f1LastInterp*(1.f-tRatio) + f1NextInterp*tRatio; - dist.f[DIR_PM0][kse] = f2LastInterp*(1.f-tRatio) + f2NextInterp*tRatio; - dist.f[DIR_P0P][kte] = f3LastInterp*(1.f-tRatio) + f3NextInterp*tRatio; - dist.f[DIR_P0M][kbe] = f4LastInterp*(1.f-tRatio) + f4NextInterp*tRatio; - dist.f[DIR_PPP][ktne] = f5LastInterp*(1.f-tRatio) + f5NextInterp*tRatio; - dist.f[DIR_PMP][ktse] = f6LastInterp*(1.f-tRatio) + f6NextInterp*tRatio; - dist.f[DIR_PPM][kbne] = f7LastInterp*(1.f-tRatio) + f7NextInterp*tRatio; - dist.f[DIR_PMM][kbse] = f8LastInterp*(1.f-tRatio) + f8NextInterp*tRatio; + dist.f[DIR_P00][ke] = f0LastInterp*(1.f-timeRatio) + f0NextInterp*timeRatio; + dist.f[DIR_PP0][kne] = f1LastInterp*(1.f-timeRatio) + f1NextInterp*timeRatio; + dist.f[DIR_PM0][kse] = f2LastInterp*(1.f-timeRatio) + f2NextInterp*timeRatio; + dist.f[DIR_P0P][kte] = f3LastInterp*(1.f-timeRatio) + f3NextInterp*timeRatio; + dist.f[DIR_P0M][kbe] = f4LastInterp*(1.f-timeRatio) + f4NextInterp*timeRatio; + dist.f[DIR_PPP][ktne] = f5LastInterp*(1.f-timeRatio) + f5NextInterp*timeRatio; + dist.f[DIR_PMP][ktse] = f6LastInterp*(1.f-timeRatio) + f6NextInterp*timeRatio; + dist.f[DIR_PPM][kbne] = f7LastInterp*(1.f-timeRatio) + f7NextInterp*timeRatio; + dist.f[DIR_PMM][kbse] = f8LastInterp*(1.f-timeRatio) + f8NextInterp*timeRatio; } //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -825,7 +825,7 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices, real* weightsSB, real* fsLast, real* fsNext, - real tRatio, + real timeRatio, unsigned long long numberOfLBnodes, bool isEvenTimestep) { @@ -955,15 +955,15 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices, getPointersToSubgridDistances(qs, subgridDistances, sizeQ); real q; - q = qs.q[DIR_P00][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P00][ke] = f0LastInterp*(1.f-tRatio) + f0NextInterp*tRatio; - q = qs.q[DIR_PP0][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PP0][kne] = f1LastInterp*(1.f-tRatio) + f1NextInterp*tRatio; - q = qs.q[DIR_PM0][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PM0][kse] = f2LastInterp*(1.f-tRatio) + f2NextInterp*tRatio; - q = qs.q[DIR_P0P][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P0P][kte] = f3LastInterp*(1.f-tRatio) + f3NextInterp*tRatio; - q = qs.q[DIR_P0M][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P0M][kbe] = f4LastInterp*(1.f-tRatio) + f4NextInterp*tRatio; - q = qs.q[DIR_PPP][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PPP][ktne] = f5LastInterp*(1.f-tRatio) + f5NextInterp*tRatio; - q = qs.q[DIR_PMP][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PMP][ktse] = f6LastInterp*(1.f-tRatio) + f6NextInterp*tRatio; - q = qs.q[DIR_PPM][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PPM][kbne] = f7LastInterp*(1.f-tRatio) + f7NextInterp*tRatio; - q = qs.q[DIR_PMM][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PMM][kbse] = f8LastInterp*(1.f-tRatio) + f8NextInterp*tRatio; + q = qs.q[DIR_P00][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P00][ke] = f0LastInterp*(1.f-timeRatio) + f0NextInterp*timeRatio; + q = qs.q[DIR_PP0][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PP0][kne] = f1LastInterp*(1.f-timeRatio) + f1NextInterp*timeRatio; + q = qs.q[DIR_PM0][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PM0][kse] = f2LastInterp*(1.f-timeRatio) + f2NextInterp*timeRatio; + q = qs.q[DIR_P0P][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P0P][kte] = f3LastInterp*(1.f-timeRatio) + f3NextInterp*timeRatio; + q = qs.q[DIR_P0M][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_P0M][kbe] = f4LastInterp*(1.f-timeRatio) + f4NextInterp*timeRatio; + q = qs.q[DIR_PPP][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PPP][ktne] = f5LastInterp*(1.f-timeRatio) + f5NextInterp*timeRatio; + q = qs.q[DIR_PMP][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PMP][ktse] = f6LastInterp*(1.f-timeRatio) + f6NextInterp*timeRatio; + q = qs.q[DIR_PPM][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PPM][kbne] = f7LastInterp*(1.f-timeRatio) + f7NextInterp*timeRatio; + q = qs.q[DIR_PMM][k]; if(q>= c0o1 && q <= c1o1) dist.f[DIR_PMM][kbse] = f8LastInterp*(1.f-timeRatio) + f8NextInterp*timeRatio; } ////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////