diff --git a/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.cpp b/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.cpp
index 8b3fca3e374f4ca08ab3a3291a49a242b0b2ff70..3026458c988c4e7624e15304f144d01f40fb0cdf 100644
--- a/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.cpp
+++ b/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.cpp
@@ -225,20 +225,20 @@ void VTKFileCollection::findFiles()
VF_LOG_CRITICAL("VTKFileCollection found no files!");
}
-void TransientBCInputFileReader::getNeighbors(uint* neighborNT, uint* neighborNB, uint* neighborST, uint* neighborSB)
+void TransientBCInputFileReader::getNeighbors(uint* neighbor0PP, uint* neighbor0PM, uint* neighbor0MP, uint* neighbor0MM)
{
- std::copy(planeNeighborNT.begin(), planeNeighborNT.end(), &neighborNT[writingOffset]);
- std::copy(planeNeighborNB.begin(), planeNeighborNB.end(), &neighborNB[writingOffset]);
- std::copy(planeNeighborST.begin(), planeNeighborST.end(), &neighborST[writingOffset]);
- std::copy(planeNeighborSB.begin(), planeNeighborSB.end(), &neighborSB[writingOffset]);
+ std::copy(planeNeighbor0PP.begin(), planeNeighbor0PP.end(), &neighbor0PP[writingOffset]);
+ std::copy(planeNeighbor0PM.begin(), planeNeighbor0PM.end(), &neighbor0PM[writingOffset]);
+ std::copy(planeNeighbor0MP.begin(), planeNeighbor0MP.end(), &neighbor0MP[writingOffset]);
+ std::copy(planeNeighbor0MM.begin(), planeNeighbor0MM.end(), &neighbor0MM[writingOffset]);
}
-void TransientBCInputFileReader::getWeights(real* _weightsNT, real* _weightsNB, real* _weightsST, real* _weightsSB)
+void TransientBCInputFileReader::getWeights(real* _weights0PP, real* _weights0PM, real* _weights0MP, real* _weights0MM)
{
- std::copy(weightsNT.begin(), weightsNT.end(), &_weightsNT[writingOffset]);
- std::copy(weightsNB.begin(), weightsNB.end(), &_weightsNB[writingOffset]);
- std::copy(weightsST.begin(), weightsST.end(), &_weightsST[writingOffset]);
- std::copy(weightsSB.begin(), weightsSB.end(), &_weightsSB[writingOffset]);
+ std::copy(weights0PP.begin(), weights0PP.end(), &_weights0PP[writingOffset]);
+ std::copy(weights0PM.begin(), weights0PM.end(), &_weights0PM[writingOffset]);
+ std::copy(weights0MP.begin(), weights0MP.end(), &_weights0MP[writingOffset]);
+ std::copy(weights0MM.begin(), weights0MM.end(), &_weights0MM[writingOffset]);
}
@@ -263,22 +263,22 @@ void VTKReader::fillArrays(std::vector<real>& coordsY, std::vector<real>& coords
real eps = 1e-7; // small number to avoid division by zero
bool perfect_match = true;
- this->weightsNT.reserve(this->nPoints);
- this->weightsNB.reserve(this->nPoints);
- this->weightsST.reserve(this->nPoints);
- this->weightsSB.reserve(this->nPoints);
+ this->weights0PP.reserve(this->nPoints);
+ this->weights0PM.reserve(this->nPoints);
+ this->weights0MP.reserve(this->nPoints);
+ this->weights0MM.reserve(this->nPoints);
- this->planeNeighborNT.reserve(this->nPoints);
- this->planeNeighborNB.reserve(this->nPoints);
- this->planeNeighborST.reserve(this->nPoints);
- this->planeNeighborSB.reserve(this->nPoints);
+ this->planeNeighbor0PP.reserve(this->nPoints);
+ this->planeNeighbor0PM.reserve(this->nPoints);
+ this->planeNeighbor0MP.reserve(this->nPoints);
+ this->planeNeighbor0MM.reserve(this->nPoints);
for(uint i=0; i<nPoints; i++)
{
real posY = coordsY[i];
real posZ = coordsZ[i];
- bool foundNT = false, foundNB = false, foundST = false, foundSB = false, foundAll = false;
+ bool found0PP = false, found0PM = false, found0MP = false, found0MM = false, foundAll = false;
uint level = this->readLevel;
@@ -289,80 +289,80 @@ void VTKReader::fillArrays(std::vector<real>& coordsY, std::vector<real>& coords
// y in simulation is x in precursor/file, z in simulation is y in precursor/file
// simulation -> file: N -> E, S -> W, T -> N, B -> S
- int idx = file.findNeighborWSB(posY, posZ, 0.f); //!> index of nearest WSB neighbor on precursor file
+ int idx = file.findNeighborMMM(posY, posZ, 0.f); //!> index of nearest WSB neighbor on precursor file
if(idx!=-1)
{
// Filter for exact matches
if(abs(posY-file.getX(idx)) < max_diff && abs(posZ-file.getY(idx)) < max_diff)
{
- this->weightsNT.emplace_back(1e6f);
- this->weightsNB.emplace_back(0.f);
- this->weightsST.emplace_back(0.f);
- this->weightsSB.emplace_back(0.f);
+ this->weights0PP.emplace_back(1e6f);
+ this->weights0PM.emplace_back(0.f);
+ this->weights0MP.emplace_back(0.f);
+ this->weights0MM.emplace_back(0.f);
uint writeIdx = this->getWriteIndex(level, fileId, idx); //!> writeIdx: index on host/device array where precursor value will be written to after loading from file
- this->planeNeighborNT.push_back(writeIdx); //!> neighbor lists mapping where BC kernel should read from on host/device array
- this->planeNeighborNB.push_back(writeIdx);
- this->planeNeighborST.push_back(writeIdx);
- this->planeNeighborSB.push_back(writeIdx);
- foundNT = true; foundNB = true; foundSB = true; foundST = true;
+ this->planeNeighbor0PP.push_back(writeIdx); //!> neighbor lists mapping where BC kernel should read from on host/device array
+ this->planeNeighbor0PM.push_back(writeIdx);
+ this->planeNeighbor0MP.push_back(writeIdx);
+ this->planeNeighbor0MM.push_back(writeIdx);
+ found0PP = true; found0PM = true; found0MM = true; found0MP = true;
}
else
{
perfect_match = false;
}
- if(!foundSB)
+ if(!found0MM)
{
- foundSB = true;
+ found0MM = true;
real dy = file.getX(idx)-posY;
real dz = file.getY(idx)-posZ;
- this->weightsSB.emplace_back(1.f/(dy*dy+dz*dz+eps));
- this->planeNeighborSB.emplace_back(getWriteIndex(level, fileId, idx));
+ this->weights0MM.emplace_back(1.f/(dy*dy+dz*dz+eps));
+ this->planeNeighbor0MM.emplace_back(getWriteIndex(level, fileId, idx));
}
}
- if(!foundNT) //NT in simulation is EN in precursor
+ if(!found0PP) //NT in simulation is EN in precursor
{
- int idx = file.findNeighborENB(posY, posZ, 0.f);
+ int idx = file.findNeighborPPM(posY, posZ, 0.f);
if(idx!=-1)
{
- foundNT = true;
+ found0PP = true;
real dy = file.getX(idx)-posY;
real dz = file.getY(idx)-posZ;
- this->weightsNT.emplace_back(1.f/(dy*dy+dz*dz+eps));
- this->planeNeighborNT.emplace_back(getWriteIndex(level, fileId, idx));
+ this->weights0PP.emplace_back(1.f/(dy*dy+dz*dz+eps));
+ this->planeNeighbor0PP.emplace_back(getWriteIndex(level, fileId, idx));
}
}
- if(!foundNB) //NB in simulation is ES in precursor
+ if(!found0PM) //NB in simulation is ES in precursor
{
- int idx = file.findNeighborESB(posY, posZ, 0.f);
+ int idx = file.findNeighborPMM(posY, posZ, 0.f);
if(idx!=-1)
{
- foundNB = true;
+ found0PM = true;
real dy = file.getX(idx)-posY;
real dz = file.getY(idx)-posZ;
- this->weightsNB.emplace_back(1.f/(dy*dy+dz*dz+eps));
- this->planeNeighborNT.emplace_back(getWriteIndex(level, fileId, idx));
+ this->weights0PM.emplace_back(1.f/(dy*dy+dz*dz+eps));
+ this->planeNeighbor0PP.emplace_back(getWriteIndex(level, fileId, idx));
}
}
- if(!foundST) //ST in simulation is WN in precursor
+ if(!found0MP) //ST in simulation is WN in precursor
{
- int idx = file.findNeighborWNB(posY, posZ, 0.f);
+ int idx = file.findNeighborMPM(posY, posZ, 0.f);
if(idx!=-1)
{
- foundST = true;
+ found0MP = true;
real dy = file.getX(idx)-posY;
real dz = file.getY(idx)-posZ;
- this->weightsST.emplace_back(1.f/(dy*dy+dz*dz+eps));
- this->planeNeighborST.emplace_back(getWriteIndex(level, fileId, idx));
+ this->weights0MP.emplace_back(1.f/(dy*dy+dz*dz+eps));
+ this->planeNeighbor0MP.emplace_back(getWriteIndex(level, fileId, idx));
}
}
- foundAll = foundNT && foundNB && foundST && foundSB;
+ foundAll = found0PP && found0PM && found0MP && found0MM;
if(foundAll) break;
}
@@ -406,35 +406,35 @@ void VTKReader::getNextData(real* data, uint numberOfNodes, real time)
uint level = this->readLevel;
for(size_t id=0; id<this->fileCollection->files[level].size(); id++)
{
- size_t nF = this->nFile[level][id];
+ size_t numberOfFiles = this->nFile[level][id];
- if(!this->fileCollection->files[level][id][nF].inZBounds(time))
+ if(!this->fileCollection->files[level][id][numberOfFiles].inZBounds(time))
{
- nF++;
+ numberOfFiles++;
- printf("switching to precursor file no. %zd\n", nF);
- if(nF == this->fileCollection->files[level][id].size())
+ printf("switching to precursor file no. %zd\n", numberOfFiles);
+ if(numberOfFiles == this->fileCollection->files[level][id].size())
throw std::runtime_error("Not enough Precursor Files to read");
- this->fileCollection->files[level][id][nF-1].unloadFile();
- if(nF+1<this->fileCollection->files[level][id].size())
+ this->fileCollection->files[level][id][numberOfFiles-1].unloadFile();
+ if(numberOfFiles+1<this->fileCollection->files[level][id].size())
{
- VTKFile* nextFile = &this->fileCollection->files[level][id][nF+1];
+ VTKFile* nextFile = &this->fileCollection->files[level][id][numberOfFiles+1];
if(! nextFile->isLoaded())
{
read.wait();
- read = std::async(std::launch::async, [](VTKFile* file){ file->loadFile(); }, &this->fileCollection->files[level][id][nF+1]);
+ read = std::async(std::launch::async, [](VTKFile* file){ file->loadFile(); }, &this->fileCollection->files[level][id][numberOfFiles+1]);
}
}
}
- VTKFile* file = &this->fileCollection->files[level][id][nF];
+ VTKFile* file = &this->fileCollection->files[level][id][numberOfFiles];
int off = file->getClosestIdxZ(time)*file->getNumberOfPointsInXYPlane();
file->getData(data, numberOfNodes, this->readIndices[level][id], this->writeIndices[level][id], off, this->writingOffset);
- this->nFile[level][id] = nF;
+ this->nFile[level][id] = numberOfFiles;
}
// }
}
\ No newline at end of file
diff --git a/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.h b/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.h
index 49b085973a39181cf0cfe23d44e818d3061ed4ca..5bee61e194670c74ca8bd8da87f3881956fff466 100644
--- a/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.h
+++ b/src/gpu/GridGenerator/TransientBCSetter/TransientBCSetter.h
@@ -46,23 +46,23 @@ public:
bool inXBounds(real posX){ return posX<=maxX && posX>=minX; };
bool inYBounds(real posY){ return posY<=maxY && posY>=minY; };
bool inZBounds(real posZ){ return posZ<=maxZ && posZ>=minZ; };
- int findNeighborWSB(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX) , getIdxSY(posY) , getIdxBZ(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborWST(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX) , getIdxSY(posY) , getIdxBZ(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborWNB(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX) , getIdxSY(posY)+1, getIdxBZ(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborWNT(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX) , getIdxSY(posY)+1, getIdxBZ(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborESB(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX)+1, getIdxSY(posY) , getIdxBZ(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborEST(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX)+1, getIdxSY(posY) , getIdxBZ(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborENB(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX)+1, getIdxSY(posY)+1, getIdxBZ(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
- int findNeighborENT(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxWX(posX)+1, getIdxSY(posY)+1, getIdxBZ(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborMMM(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX) , getIdx0M0(posY) , getIdx00M(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborMMP(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX) , getIdx0M0(posY) , getIdx00M(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborMPM(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX) , getIdx0M0(posY)+1, getIdx00M(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborMPP(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX) , getIdx0M0(posY)+1, getIdx00M(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborPMM(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX)+1, getIdx0M0(posY) , getIdx00M(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborPMP(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX)+1, getIdx0M0(posY) , getIdx00M(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborPPM(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX)+1, getIdx0M0(posY)+1, getIdx00M(posZ) ); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
+ int findNeighborPPP(real posX, real posY, real posZ){ int idx = getLinearIndex(getIdxM00(posX)+1, getIdx0M0(posY)+1, getIdx00M(posZ)+1); return (idx>=0) && (idx<nx*ny*nz) ? idx : -1; };
int getIdxX(int linearIdx){ return linearIdx%nx;};
int getIdxY(int linearIdx){ return (linearIdx/nx)%ny;};
int getIdxZ(int linearIdx){ return linearIdx/(nx*ny); };
real getX(int linearIdx){ return getIdxX(linearIdx)*deltaX+minX; };
real getY(int linearIdx){ return getIdxY(linearIdx)*deltaY+minY; };
real getZ(int linearIdx){ return getIdxZ(linearIdx)*deltaZ+minZ; };
- int getIdxWX(real posX){ return (posX-minX)/deltaX; };
- int getIdxSY(real posY){ return (posY-minY)/deltaY; };
- int getIdxBZ(real posZ){ return (posZ-minZ)/deltaZ; };
+ int getIdxM00(real posX){ return (posX-minX)/deltaX; };
+ int getIdx0M0(real posY){ return (posY-minY)/deltaY; };
+ int getIdx00M(real posZ){ return (posZ-minZ)/deltaZ; };
int getClosestIdxX(real posX){ int x = round((posX-minX)/deltaX); return x>nx ? nx : (x<0 ? 0 : x);};
int getClosestIdxY(real posY){ int y = round((posY-minY)/deltaY); return y>ny ? ny : (y<0 ? 0 : y);};
int getClosestIdxZ(real posZ){ int z = round((posZ-minZ)/deltaZ); return z>nz ? nz : (z<0 ? 0 : z);};
@@ -156,12 +156,12 @@ public:
uint getNPointsRead(){return nPointsRead; };
size_t getNumberOfQuantities(){ return nQuantities; };
void setWritingOffset(uint offset){ this->writingOffset = offset; }
- void getNeighbors(uint* neighborNT, uint* neighborNB, uint* neighborST, uint* neighborSN);
- void getWeights(real* _weightsNT, real* _weightsNB, real* _weightsST, real* _weightsSB);
+ void getNeighbors(uint* neighbor0PP, uint* neighbor0PM, uint* neighbor0MP, uint* neighbor0MM);
+ void getWeights(real* _weights0PP, real* _weights0PM, real* _weights0MP, real* _weights0MM);
public:
- std::vector<uint> planeNeighborNT, planeNeighborNB, planeNeighborST, planeNeighborSB;
- std::vector<real> weightsNT, weightsNB, weightsST, weightsSB;
+ std::vector<uint> planeNeighbor0PP, planeNeighbor0PM, planeNeighbor0MP, planeNeighbor0MM;
+ std::vector<real> weights0PP, weights0PM, weights0MP, weights0MM;
protected:
uint nPoints, nPointsRead, writingOffset;
diff --git a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
index 9b430188235845d416e085f55a4586ee4b7c8f20..0900c2d587ba9811c480427b833e7e083216cf10 100644
--- a/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
+++ b/src/gpu/GridGenerator/grid/BoundaryConditions/Side.cpp
@@ -73,8 +73,7 @@ void Side::addIndices(SPtr<Grid> grid, SPtr<BoundaryCondition> boundaryCondition
|| grid->getFieldEntry(index) == vf::gpu::BC_VELOCITY
|| grid->getFieldEntry(index) == vf::gpu::BC_NOSLIP
|| grid->getFieldEntry(index) == vf::gpu::BC_SLIP
- || grid->getFieldEntry(index) == vf::gpu::BC_STRESS )
- /*&& boundaryCondition->getType()!=vf::gpu::BC_STRESS*/ )
+ || grid->getFieldEntry(index) == vf::gpu::BC_STRESS ))
{
grid->setFieldEntry(index, boundaryCondition->getType());
boundaryCondition->indices.push_back(index);
@@ -177,13 +176,6 @@ void Side::setQs(SPtr<Grid> grid, SPtr<BoundaryCondition> boundaryCondition, uin
this->getNormal()[1]*grid->getDirection()[dir * DIMENSION + 1]+
this->getNormal()[2]*grid->getDirection()[dir * DIMENSION + 2] ) > 0;
- // if(boundaryCondition->getType()==vf::gpu::BC_VELOCITY && z < 8.0 )
- // {
- // alignedWithNormal = true;
- // printf("XYZ: %f \t %f \t %f \n", x,y,z);
- // printf("dir: %d \t %d \t %d \n\n", grid->getDirection()[dir * DIMENSION + 0], grid->getDirection()[dir * DIMENSION + 1], grid->getDirection()[dir * DIMENSION + 2]);
- // }
-
uint neighborIndex = grid->transCoordToIndex( neighborX, neighborY, neighborZ );
if((grid->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID_BOUNDARY ||
grid->getFieldEntry(neighborIndex) == vf::gpu::STOPPER_OUT_OF_GRID ||
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
index 8c3db0afdf89ceb460b87d6b0b0ccd91db56f428..065665d216e3cf7904530b94c8bb6480bb565c8a 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/GridBuilder.h
@@ -115,8 +115,8 @@ public:
virtual void getPressureQs(real *qs[27], int level) const = 0;
virtual uint getPrecursorSize(int level) const = 0;
- virtual void getPrecursorValues(uint* neighborNT, uint* neighborNB, uint* neighborST, uint* neighborSB,
- real* weightsNT, real* weightsNB, real* weightsST, real* weightsSB,
+ virtual void getPrecursorValues(uint* neighbor0PP, uint* neighbor0PM, uint* neighbor0MP, uint* neighbor0MM,
+ real* weights0PP, real* weights0PM, real* weights0MP, real* weights0MM,
int* indices, std::vector<SPtr<TransientBCInputFileReader>>& reader,
int& numberOfPrecursorNodes, size_t& numberOfQuantities, uint& timeStepsBetweenReads,
real& velocityX, real& velocityY, real& velocityZ, int level) const = 0;
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
index ec0db541b60c8b278ff3926f8a6fd0cad540b1d3..ff6f4913e4cd8a32c05272ef583f90a2cf226edc 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp
@@ -644,8 +644,8 @@ uint LevelGridBuilder::getPrecursorSize(int level) const
return size;
}
-void LevelGridBuilder::getPrecursorValues( uint* neighborNT, uint* neighborNB, uint* neighborST, uint* neighborSB,
- real* weightsNT, real* weightsNB, real* weightsST, real* weightsSB,
+void LevelGridBuilder::getPrecursorValues( uint* neighbor0PP, uint* neighbor0PM, uint* neighbor0MP, uint* neighbor0MM,
+ real* weights0PP, real* weights0PM, real* weights0MP, real* weights0MM,
int* indices, std::vector<SPtr<TransientBCInputFileReader>>& reader,
int& numberOfPrecursorNodes, size_t& numberOfQuantities, uint& timeStepsBetweenReads,
real& velocityX, real& velocityY, real& velocityZ, int level) const
@@ -653,7 +653,7 @@ void LevelGridBuilder::getPrecursorValues( uint* neighborNT, uint* neighborNB,
int allIndicesCounter = 0;
int allNodesCounter = 0;
uint tmpTimeStepsBetweenReads = 0;
- size_t tmpNQuantities = 0;
+ size_t tmpNumberOfQuantities = 0;
for (auto boundaryCondition : boundaryConditions[level]->precursorBoundaryConditions)
{
@@ -676,11 +676,11 @@ void LevelGridBuilder::getPrecursorValues( uint* neighborNT, uint* neighborNB,
allIndicesCounter++;
}
BCreader->fillArrays(y, z);
- BCreader->getNeighbors(neighborNT, neighborNB, neighborST, neighborSB);
- BCreader->getWeights(weightsNT, weightsNB, weightsST, weightsSB);
- if(tmpNQuantities == 0)
- tmpNQuantities = BCreader->getNumberOfQuantities();
- if(tmpNQuantities != BCreader->getNumberOfQuantities())
+ BCreader->getNeighbors(neighbor0PP, neighbor0PM, neighbor0MP, neighbor0MM);
+ BCreader->getWeights(weights0PP, weights0PM, weights0MP, weights0MM);
+ if(tmpNumberOfQuantities == 0)
+ tmpNumberOfQuantities = BCreader->getNumberOfQuantities();
+ if(tmpNumberOfQuantities != BCreader->getNumberOfQuantities())
throw std::runtime_error("All precursor files must have the same quantities.");
allNodesCounter += BCreader->getNPointsRead();
velocityX = boundaryCondition->getVelocityX();
@@ -693,9 +693,9 @@ void LevelGridBuilder::getPrecursorValues( uint* neighborNT, uint* neighborNB,
throw std::runtime_error("timeStepsBetweenReads of precursor needs to be larger than 0.");
timeStepsBetweenReads = tmpTimeStepsBetweenReads;
- if (tmpNQuantities == 0)
+ if (tmpNumberOfQuantities == 0)
throw std::runtime_error("Number of quantities in precursor needs to be larger than 0.");
- numberOfQuantities = tmpNQuantities;
+ numberOfQuantities = tmpNumberOfQuantities;
}
void LevelGridBuilder::getPrecursorQs(real* qs[27], int level) const
diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
index afaf4118469a3cc935cd862d311629fabd1c0379..56ae1e4fce6185591fba97f49ba504ced259aea5 100644
--- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
+++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.h
@@ -38,6 +38,8 @@
#include <memory>
#include <array>
+#include <lbm/constants/NumericConstants.h>
+
#include "gpu/GridGenerator/global.h"
#include "gpu/GridGenerator/grid/GridBuilder/GridBuilder.h"
@@ -45,6 +47,8 @@
#include "gpu/GridGenerator/grid/GridInterface.h"
#include "gpu/GridGenerator/grid/NodeValues.h"
+using namespace vf::lbm::constant;
+
struct Vertex;
class Grid;
class Transformator;
@@ -64,8 +68,6 @@ enum class SideType;
class TransientBCInputFileReader;
class FileCollection;
-
-
class LevelGridBuilder : public GridBuilder
{
protected:
@@ -85,7 +87,7 @@ public:
GRIDGENERATOR_EXPORT void setPeriodicBoundaryCondition(bool periodic_X, bool periodic_Y, bool periodic_Z);
GRIDGENERATOR_EXPORT void setNoSlipBoundaryCondition(SideType sideType);
GRIDGENERATOR_EXPORT void setPrecursorBoundaryCondition(SideType sideType, SPtr<FileCollection> fileCollection, int timeStepsBetweenReads,
- real velocityX=0.0f, real velocityY=0.0f, real velocityZ=0.0f,
+ real velocityX=c0o1, real velocityY=c0o1, real velocityZ=c0o1,
std::vector<uint> fileLevelToGridLevelMap = {});
GRIDGENERATOR_EXPORT void setEnableFixRefinementIntoTheWall(bool enableFixRefinementIntoTheWall);
@@ -129,8 +131,8 @@ public:
GRIDGENERATOR_EXPORT virtual void getPressureQs(real* qs[27], int level) const override;
GRIDGENERATOR_EXPORT uint getPrecursorSize(int level) const override;
- GRIDGENERATOR_EXPORT void getPrecursorValues( uint* neighborNT, uint* neighborNB, uint* neighborST, uint* neighborSB,
- real* weightsNT, real* weightsNB, real* weightsST, real* weightsSB,
+ GRIDGENERATOR_EXPORT void getPrecursorValues( uint* neighbor0PP, uint* neighbor0PM, uint* neighbor0MP, uint* neighbor0MM,
+ real* weights0PP, real* weights0PM, real* weights0MP, real* weights0MM,
int* indices, std::vector<SPtr<TransientBCInputFileReader>>& reader,
int& numberOfPrecursorNodes, size_t& numberOfQuantities, uint& timeStepsBetweenReads,
real& velocityX, real& velocityY, real& velocityZ, int level) const override;
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index c3dd9e41495645d06e54514cc25eb336ed6fc9c5..d89f4cd85985694b4529d338ab5e4c5a86cd3a38 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -342,10 +342,10 @@ void GridGenerator::allocArrays_BoundaryValues()
cudaMemoryManager->cudaAllocPrecursorBC(level);
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
builder->getPrecursorValues(
- para->getParH(level)->precursorBC.planeNeighborNT, para->getParH(level)->precursorBC.planeNeighborNB,
- para->getParH(level)->precursorBC.planeNeighborST, para->getParH(level)->precursorBC.planeNeighborSB,
- para->getParH(level)->precursorBC.weightsNT, para->getParH(level)->precursorBC.weightsNB,
- para->getParH(level)->precursorBC.weightsST, para->getParH(level)->precursorBC.weightsSB,
+ para->getParH(level)->precursorBC.planeNeighbor0PP, para->getParH(level)->precursorBC.planeNeighbor0PM,
+ para->getParH(level)->precursorBC.planeNeighbor0MP, para->getParH(level)->precursorBC.planeNeighbor0MM,
+ para->getParH(level)->precursorBC.weights0PP, para->getParH(level)->precursorBC.weights0PM,
+ para->getParH(level)->precursorBC.weights0MP, para->getParH(level)->precursorBC.weights0MM,
para->getParH(level)->precursorBC.k, para->getParH(level)->transientBCInputFileReader, para->getParH(level)->precursorBC.numberOfPrecursorNodes,
para->getParH(level)->precursorBC.numberOfQuantities, para->getParH(level)->precursorBC.timeStepsBetweenReads,
para->getParH(level)->precursorBC.velocityX, para->getParH(level)->precursorBC.velocityY, para->getParH(level)->precursorBC.velocityZ,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 8807c22300af92b294976c3d59317f4118e0b72e..c4f3807efe513d7f8d17d9cd268d55febee47a00 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -1669,28 +1669,28 @@ void CudaMemoryManager::cudaAllocPrecursorBC(int lev)
checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.q27[0], parameter->getD3Qxx()*memSizeQReal));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighborNT, memSizeQUint));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighborNB, memSizeQUint));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighborST, memSizeQUint));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighborSB, memSizeQUint));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighbor0PP, memSizeQUint));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighbor0PM, memSizeQUint));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighbor0MP, memSizeQUint));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.planeNeighbor0MM, memSizeQUint));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weightsNT, memSizeQReal));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weightsNB, memSizeQReal));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weightsST, memSizeQReal));
- checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weightsSB, memSizeQReal));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weights0PP, memSizeQReal));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weights0PM, memSizeQReal));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weights0MP, memSizeQReal));
+ checkCudaErrors( cudaMallocHost((void**) ¶meter->getParH(lev)->precursorBC.weights0MM, memSizeQReal));
checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.k, memSizeQInt));
checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.q27[0], parameter->getD3Qxx()*memSizeQReal));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighborNT, memSizeQUint));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighborNB, memSizeQUint));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighborST, memSizeQUint));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighborSB, memSizeQUint));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighbor0PP, memSizeQUint));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighbor0PM, memSizeQUint));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighbor0MP, memSizeQUint));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.planeNeighbor0MM, memSizeQUint));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weightsNT, memSizeQReal));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weightsNB, memSizeQReal));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weightsST, memSizeQReal));
- checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weightsSB, memSizeQReal));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weights0PP, memSizeQReal));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weights0PM, memSizeQReal));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weights0MP, memSizeQReal));
+ checkCudaErrors( cudaMalloc((void**) ¶meter->getParD(lev)->precursorBC.weights0MM, memSizeQReal));
real memSize = memSizeQInt+4*memSizeQUint+(4+parameter->getD3Qxx())*memSizeQReal;
setMemsizeGPU(memSize, false);
@@ -1723,15 +1723,15 @@ void CudaMemoryManager::cudaCopyPrecursorBC(int lev)
checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.q27[0], parameter->getParH(lev)->precursorBC.q27[0], memSizeQReal*parameter->getD3Qxx(), cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighborNT, parameter->getParH(lev)->precursorBC.planeNeighborNT, memSizeQUint, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighborNB, parameter->getParH(lev)->precursorBC.planeNeighborNB, memSizeQUint, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighborST, parameter->getParH(lev)->precursorBC.planeNeighborST, memSizeQUint, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighborSB, parameter->getParH(lev)->precursorBC.planeNeighborSB, memSizeQUint, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighbor0PP, parameter->getParH(lev)->precursorBC.planeNeighbor0PP, memSizeQUint, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighbor0PM, parameter->getParH(lev)->precursorBC.planeNeighbor0PM, memSizeQUint, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighbor0MP, parameter->getParH(lev)->precursorBC.planeNeighbor0MP, memSizeQUint, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.planeNeighbor0MM, parameter->getParH(lev)->precursorBC.planeNeighbor0MM, memSizeQUint, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weightsNT, parameter->getParH(lev)->precursorBC.weightsNT, memSizeQReal, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weightsNB, parameter->getParH(lev)->precursorBC.weightsNB, memSizeQReal, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weightsST, parameter->getParH(lev)->precursorBC.weightsST, memSizeQReal, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weightsSB, parameter->getParH(lev)->precursorBC.weightsSB, memSizeQReal, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weights0PP, parameter->getParH(lev)->precursorBC.weights0PP, memSizeQReal, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weights0PM, parameter->getParH(lev)->precursorBC.weights0PM, memSizeQReal, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weights0MP, parameter->getParH(lev)->precursorBC.weights0MP, memSizeQReal, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->precursorBC.weights0MM, parameter->getParH(lev)->precursorBC.weights0MM, memSizeQReal, cudaMemcpyHostToDevice));
}
void CudaMemoryManager::cudaCopyPrecursorData(int lev)
{
@@ -1749,29 +1749,29 @@ void CudaMemoryManager::cudaFreePrecursorBC(int lev)
checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.q27[0]));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighborNT));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighborNB));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighborST));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighborSB));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighbor0PP));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighbor0PM));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighbor0MP));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.planeNeighbor0MM));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weightsNT));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weightsNB));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weightsST));
- checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weightsSB));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weights0PP));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weights0PM));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weights0MP));
+ checkCudaErrors( cudaFreeHost( parameter->getParH(lev)->precursorBC.weights0MM));
checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.k));
checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.q27[0]));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighborNT));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighborNB));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighborST));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighborSB));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighbor0PP));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighbor0PM));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighbor0MP));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.planeNeighbor0MM));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weightsNT));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weightsNB));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weightsST));
- checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weightsSB));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weights0PP));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weights0PM));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weights0MP));
+ checkCudaErrors( cudaFree( parameter->getParD(lev)->precursorBC.weights0MM));
}
void CudaMemoryManager::cudaFreePrecursorData(int lev)
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
index 9ef5057bbf12af887d438e49b974402136fc60c1..a7b22c244da368b9591c5f528bb7bd5eaea91e50 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
@@ -1257,10 +1257,10 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* vLast,
real* vCurrent,
real velocityX,
@@ -1283,10 +1283,10 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* vLast,
real* vCurrent,
real velocityX,
@@ -1308,10 +1308,10 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* fsLast,
real* fsNext,
real timeRatio,
@@ -1330,10 +1330,10 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* fsLast,
real* fsNext,
real timeRatio,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index 06628edb5e3d47f5b0ed44ce61ec6da134c551da..41f74113fd90b76209887f5624c60f2ca84ad533 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -3216,8 +3216,8 @@ void QPrecursorDevCompZeroPress(LBMSimulationParameter* parameterDevice, QforPre
QPrecursorDeviceCompZeroPress<<< grid.grid, grid.threads >>>(boundaryCondition->k, boundaryCondition->numberOfBCnodes, boundaryCondition->numberOfPrecursorNodes, boundaryCondition->sizeQ, parameterDevice->omega,
parameterDevice->distributions.f[0], boundaryCondition->q27[0],
parameterDevice->neighborX, parameterDevice->neighborY, parameterDevice->neighborZ,
- boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB,
- boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB,
+ boundaryCondition->planeNeighbor0PP, boundaryCondition->planeNeighbor0PM, boundaryCondition->planeNeighbor0MP, boundaryCondition->planeNeighbor0MM,
+ boundaryCondition->weights0PP, boundaryCondition->weights0PM, boundaryCondition->weights0MP, boundaryCondition->weights0MM,
boundaryCondition->last, boundaryCondition->current,
boundaryCondition->velocityX, boundaryCondition->velocityY, boundaryCondition->velocityZ,
timeRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep);
@@ -3232,8 +3232,8 @@ void PrecursorDevEQ27( LBMSimulationParameter* parameterDevice, QforPrecursorBou
PrecursorDeviceEQ27<<< grid.grid, grid.threads >>>(boundaryCondition->k, boundaryCondition->numberOfBCnodes, boundaryCondition->numberOfPrecursorNodes, parameterDevice->omega, parameterDevice->distributions.f[0],
parameterDevice->neighborX, parameterDevice->neighborX, parameterDevice->neighborX,
- boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB,
- boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB,
+ boundaryCondition->planeNeighbor0PP, boundaryCondition->planeNeighbor0PM, boundaryCondition->planeNeighbor0MP, boundaryCondition->planeNeighbor0MM,
+ boundaryCondition->weights0PP, boundaryCondition->weights0PM, boundaryCondition->weights0MP, boundaryCondition->weights0MM,
boundaryCondition->last, boundaryCondition->current,
boundaryCondition->velocityX, boundaryCondition->velocityY, boundaryCondition->velocityZ,
timeRatio, velocityRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep);
@@ -3248,8 +3248,8 @@ void PrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPre
PrecursorDeviceDistributions<<< grid.grid, grid.threads >>>(boundaryCondition->k, boundaryCondition->numberOfBCnodes, boundaryCondition->numberOfPrecursorNodes, parameterDevice->distributions.f[0],
parameterDevice->neighborX, parameterDevice->neighborY, parameterDevice->neighborZ,
- boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB,
- boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB,
+ boundaryCondition->planeNeighbor0PP, boundaryCondition->planeNeighbor0PM, boundaryCondition->planeNeighbor0MP, boundaryCondition->planeNeighbor0MM,
+ boundaryCondition->weights0PP, boundaryCondition->weights0PM, boundaryCondition->weights0MP, boundaryCondition->weights0MM,
boundaryCondition->last, boundaryCondition->current,
timeRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep);
getLastCudaError("QPrecursorDeviceCompZeroPress execution failed");
@@ -3264,8 +3264,8 @@ void QPrecursorDevDistributions( LBMSimulationParameter* parameterDevice, QforPr
QPrecursorDeviceDistributions<<< grid.grid, grid.threads >>>(boundaryCondition->k, boundaryCondition->q27[0], boundaryCondition->sizeQ, boundaryCondition->numberOfBCnodes, boundaryCondition->numberOfPrecursorNodes, parameterDevice->distributions.f[0],
parameterDevice->neighborX, parameterDevice->neighborY, parameterDevice->neighborZ,
- boundaryCondition->planeNeighborNT, boundaryCondition->planeNeighborNB, boundaryCondition->planeNeighborST, boundaryCondition->planeNeighborSB,
- boundaryCondition->weightsNT, boundaryCondition->weightsNB, boundaryCondition->weightsST, boundaryCondition->weightsSB,
+ boundaryCondition->planeNeighbor0PP, boundaryCondition->planeNeighbor0PM, boundaryCondition->planeNeighbor0MP, boundaryCondition->planeNeighbor0MM,
+ boundaryCondition->weights0PP, boundaryCondition->weights0PM, boundaryCondition->weights0MP, boundaryCondition->weights0MM,
boundaryCondition->last, boundaryCondition->current,
timeRatio, parameterDevice->numberOfNodes, parameterDevice->isEvenTimestep);
getLastCudaError("QPrecursorDeviceCompZeroPress execution failed");
diff --git a/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu b/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu
index f089a32955af615bfa5744857e1c7d76e2bb42d3..e7b0c60797a64e8068e9abd2d9e8e0c1e577c3c6 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/PrecursorBCs27.cu
@@ -23,10 +23,10 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* vLast,
real* vCurrent,
real velocityX,
@@ -47,7 +47,7 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices,
real vxNextInterpd, vyNextInterpd, vzNextInterpd;
uint kNT = neighborsNT[k];
- real dNT = weightsNT[k];
+ real dNT = weights0PP[k];
real* vxLast = vLast;
real* vyLast = &vLast[numberOfPrecursorNodes];
@@ -63,9 +63,9 @@ __global__ void QPrecursorDeviceCompZeroPress( int* subgridDistanceIndices,
uint kST = neighborsST[k];
uint kSB = neighborsSB[k];
- real dNB = weightsNB[k];
- real dST = weightsST[k];
- real dSB = weightsSB[k];
+ real dNB = weights0PM[k];
+ real dST = weights0MP[k];
+ real dSB = weights0MM[k];
real invWeightSum = 1.f/(dNT+dNB+dST+dSB);
@@ -436,10 +436,10 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* vLast,
real* vCurrent,
real velocityX,
@@ -460,7 +460,7 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices,
real vxNextInterpd, vyNextInterpd, vzNextInterpd;
uint kNT = neighborsNT[k];
- real dNT = weightsNT[k];
+ real dNT = weights0PP[k];
real* vxLast = vLast;
real* vyLast = &vLast[numberOfPrecursorNodes];
@@ -476,9 +476,9 @@ __global__ void PrecursorDeviceEQ27( int* subgridDistanceIndices,
uint kST = neighborsST[k];
uint kSB = neighborsSB[k];
- real dNB = weightsNB[k];
- real dST = weightsST[k];
- real dSB = weightsSB[k];
+ real dNB = weights0PM[k];
+ real dST = weights0MP[k];
+ real dSB = weights0MM[k];
real invWeightSum = 1.f/(dNT+dNB+dST+dSB);
@@ -660,10 +660,10 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* fsLast,
real* fsNext,
real timeRatio,
@@ -675,7 +675,7 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices,
if(k>=numberOfBCnodes) return;
uint kNT = neighborsNT[k];
- real dNT = weightsNT[k];
+ real dNT = weights0PP[k];
real f0LastInterp, f1LastInterp, f2LastInterp, f3LastInterp, f4LastInterp, f5LastInterp, f6LastInterp, f7LastInterp, f8LastInterp;
real f0NextInterp, f1NextInterp, f2NextInterp, f3NextInterp, f4NextInterp, f5NextInterp, f6NextInterp, f7NextInterp, f8NextInterp;
@@ -707,9 +707,9 @@ __global__ void PrecursorDeviceDistributions( int* subgridDistanceIndices,
uint kST = neighborsST[k];
uint kSB = neighborsSB[k];
- real dNB = weightsNB[k];
- real dST = weightsST[k];
- real dSB = weightsSB[k];
+ real dNB = weights0PM[k];
+ real dST = weights0MP[k];
+ real dSB = weights0MM[k];
real invWeightSum = 1.f/(dNT+dNB+dST+dSB);
@@ -819,10 +819,10 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices,
uint* neighborsNB,
uint* neighborsST,
uint* neighborsSB,
- real* weightsNT,
- real* weightsNB,
- real* weightsST,
- real* weightsSB,
+ real* weights0PP,
+ real* weights0PM,
+ real* weights0MP,
+ real* weights0MM,
real* fsLast,
real* fsNext,
real timeRatio,
@@ -834,7 +834,7 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices,
if(k>=numberOfBCnodes) return;
uint kNT = neighborsNT[k];
- real dNT = weightsNT[k];
+ real dNT = weights0PP[k];
real f0LastInterp, f1LastInterp, f2LastInterp, f3LastInterp, f4LastInterp, f5LastInterp, f6LastInterp, f7LastInterp, f8LastInterp;
real f0NextInterp, f1NextInterp, f2NextInterp, f3NextInterp, f4NextInterp, f5NextInterp, f6NextInterp, f7NextInterp, f8NextInterp;
@@ -866,9 +866,9 @@ __global__ void QPrecursorDeviceDistributions( int* subgridDistanceIndices,
uint kST = neighborsST[k];
uint kSB = neighborsSB[k];
- real dNB = weightsNB[k];
- real dST = weightsST[k];
- real dSB = weightsSB[k];
+ real dNB = weights0PM[k];
+ real dST = weights0MP[k];
+ real dSB = weights0MM[k];
real invWeightSum = 1.f/(dNT+dNB+dST+dSB);
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index c07769d5c06f4443335564511d449e82462f046f..a7517ca0f07b9c1379501aa3ff750637d230a7bd 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -240,8 +240,8 @@ typedef struct QforPrecursorBC{
uint timeStepsBetweenReads;
size_t numberOfQuantities;
real* q27[27];
- uint* planeNeighborNT, *planeNeighborNB, *planeNeighborST, *planeNeighborSB;
- real* weightsNT, *weightsNB, *weightsST, *weightsSB;
+ uint* planeNeighbor0PP, *planeNeighbor0PM, *planeNeighbor0MP, *planeNeighbor0MM;
+ real* weights0PP, *weights0PM, *weights0MP, *weights0MM;
real* last, *current, *next;
real velocityX, velocityY, velocityZ;
}QforPrecursorBoundaryConditions;