From ab7658cf2bb11d3448d97574d1a06ba858b3c1fc Mon Sep 17 00:00:00 2001 From: Soeren Peters <peters@irmb.tu-bs.de> Date: Wed, 24 Feb 2021 11:29:05 +0100 Subject: [PATCH] Fix various msvc compiler warnings. With msvc the compiler flags are explicit passed to nvcc with -Xcompiler. --- CMake/CMakeSetCompilerFlags.cmake | 3 + apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp | 6 +- src/gpu/GridGenerator/CMakeLists.txt | 5 + .../StreetPointFinder/StreetPointFinder.cpp | 2 +- .../TriangularMesh/TriangularMesh.cu | 12 +- .../grid/GridBuilder/LevelGridBuilder.cpp | 8 +- .../grid/GridBuilder/MultipleGridBuilder.cpp | 8 +- src/gpu/GridGenerator/grid/GridImp.cu | 8 +- src/gpu/GridGenerator/grid/GridImp.h | 2 +- src/gpu/GridGenerator/grid/GridInterface.cu | 4 +- .../GridCpuStrategy/GridCpuStrategy.cpp | 8 +- .../io/STLReaderWriter/STLReader.cpp | 2 +- .../io/STLReaderWriter/STLWriter.cpp | 2 +- .../SimulationFileWriter.cpp | 2 +- .../GridReaderFiles/GridReader.cpp | 12 +- .../GridReaderGenerator/GridGenerator.cpp | 16 +- .../VirtualFluids_GPU/GPU/AdvecDiffBCs27.cu | 1149 +++++++++-------- .../VirtualFluids_GPU/GPU/Cumulant_F3_27.cu | 48 +- src/gpu/VirtualFluids_GPU/GPU/Particles.cu | 733 +++++------ src/gpu/VirtualFluids_GPU/GPU/ScaleFC27.cu | 6 +- .../VirtualFluids_GPU/GPU/ScaleFC_F3_27.cu | 4 +- .../Input/PositionReader.cpp | 2 +- .../CumulantAll4CompSP27_Device.cu | 4 +- .../Mod27/ADComp27/ADComp27_Device.cu | 2 +- .../Mod7/ADIncomp7/ADIncomp7_Device.cu | 6 +- .../PMCumulantOneCompSP27_Device.cu | 2 +- src/gpu/VirtualFluids_GPU/Output/DataWriter.h | 2 - .../VirtualFluids_GPU/Output/LogWriter.hpp | 4 +- .../VirtualFluids_GPU/Parameter/Parameter.cpp | 6 +- 29 files changed, 1039 insertions(+), 1029 deletions(-) diff --git a/CMake/CMakeSetCompilerFlags.cmake b/CMake/CMakeSetCompilerFlags.cmake index 784f3f24a..2ea8c0b2f 100644 --- a/CMake/CMakeSetCompilerFlags.cmake +++ b/CMake/CMakeSetCompilerFlags.cmake @@ -68,6 +68,9 @@ function(addAdditionalFlags project_name) # compile options foreach(flag IN LISTS CS_COMPILER_FLAGS_CXX) target_compile_options(${project_name} PRIVATE "$<$<COMPILE_LANGUAGE:CXX>:${flag}>") + if(MSVC) + target_compile_options(${project_name} PRIVATE "$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=${flag}>") + endif() endforeach() foreach(flag IN LISTS CS_COMPILER_FLAGS_CXX_DEBUG) diff --git a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp index 19cb63009..72942b001 100644 --- a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp +++ b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp @@ -93,7 +93,7 @@ const real Re = 500.0;// 1000.0; const real velocity = 1.0; -const real dt = 1.0e-3; //0.5e-3; +const real dt = (real)1.0e-3; //0.5e-3; const uint nx = 64; @@ -162,8 +162,8 @@ void multipleLevel(const std::string& configPath) const real velocityLB = velocity * dt / dx; // LB units - const real vx = velocityLB / sqrt(2.0); // LB units - const real vy = velocityLB / sqrt(2.0); // LB units + const real vx = velocityLB / (real)sqrt(2.0); // LB units + const real vy = velocityLB / (real)sqrt(2.0); // LB units const real viscosityLB = nx * velocityLB / Re; // LB units diff --git a/src/gpu/GridGenerator/CMakeLists.txt b/src/gpu/GridGenerator/CMakeLists.txt index feb01cadc..30d705e0e 100644 --- a/src/gpu/GridGenerator/CMakeLists.txt +++ b/src/gpu/GridGenerator/CMakeLists.txt @@ -22,5 +22,10 @@ if(NOT MSVC) set(cuda_warnings_suppressions "${cuda_warnings_suppressions} --diag_suppress=3311") endif() +# suppress warning 2979: calling a __host__ function from __host__ __device__ is not allowed +if(MSVC) + set(cuda_warnings_suppressions "${cuda_warnings_suppressions} --diag_suppress=2979 --diag_suppress=2976 --diag_suppress=3005") +endif() + target_compile_options(${library_name} PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Xcudafe ${cuda_warnings_suppressions} >) diff --git a/src/gpu/GridGenerator/StreetPointFinder/StreetPointFinder.cpp b/src/gpu/GridGenerator/StreetPointFinder/StreetPointFinder.cpp index 5ebad0daa..23d7c92ab 100644 --- a/src/gpu/GridGenerator/StreetPointFinder/StreetPointFinder.cpp +++ b/src/gpu/GridGenerator/StreetPointFinder/StreetPointFinder.cpp @@ -112,7 +112,7 @@ void StreetPointFinder::prepareSimulationFileData() ////////////////////////////////////////////////////////////////////////// // prepare vectors - uint numberOfCells = this->sparseIndicesLB.size(); + uint numberOfCells = (uint)this->sparseIndicesLB.size(); mapNashToConc.resize(numberOfCells); diff --git a/src/gpu/GridGenerator/geometries/TriangularMesh/TriangularMesh.cu b/src/gpu/GridGenerator/geometries/TriangularMesh/TriangularMesh.cu index a85661387..4fcf93ef4 100644 --- a/src/gpu/GridGenerator/geometries/TriangularMesh/TriangularMesh.cu +++ b/src/gpu/GridGenerator/geometries/TriangularMesh/TriangularMesh.cu @@ -60,7 +60,7 @@ Object* TriangularMesh::clone() const uint TriangularMesh::getNumberOfTriangles() const { - return triangleVec.size(); + return (uint)triangleVec.size(); } @@ -94,7 +94,7 @@ void TriangularMesh::initalizeDataFromTriangles() this->triangles = triangleVec.data(); this->size = long(triangleVec.size()); - for (uint i = 0; i < this->size; i++) { + for (std::size_t i = 0; i < this->size; i++) { this->minmax.setMinMax(this->triangleVec[i]); } } @@ -167,7 +167,7 @@ void TriangularMesh::scale(double offset) auto averrageNormals = getAverrageNormalsPerVertex(trianglesPerVertex); - for (int vertexID = 0; vertexID < this->getNumberOfTriangles() * 3; vertexID++) + for (std::size_t vertexID = 0; vertexID < this->getNumberOfTriangles() * 3; vertexID++) { int coordinatedID = finder.sortedToTriangles[vertexID][IDS::coordinateID]; Vertex averrageNormal = averrageNormals[coordinatedID]; @@ -205,13 +205,13 @@ void TriangularMesh::scale(double offset) //printf("\n\n"); averrageNormal.normalize(); - const int triangleID = vertexID / 3; - const int vertexTriangleID = vertexID % 3; + const int triangleID = (int)vertexID / 3; + const int vertexTriangleID = (int)vertexID % 3; Vertex intersection; Vertex p = this->triangleVec[triangleID].v1 + this->triangleVec[triangleID].normal * offset; Vertex lineOrigin = this->triangleVec[triangleID].get(vertexTriangleID); - bool b = intersectPlane(this->triangleVec[triangleID].normal, p, lineOrigin, averrageNormal, intersection); + //bool b = intersectPlane(this->triangleVec[triangleID].normal, p, lineOrigin, averrageNormal, intersection); triangles[triangleID].set(vertexTriangleID, intersection); triangles[triangleID].calcNormal(); diff --git a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp index 700ff5c13..2312b7198 100644 --- a/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp +++ b/src/gpu/GridGenerator/grid/GridBuilder/LevelGridBuilder.cpp @@ -315,9 +315,9 @@ void LevelGridBuilder::getVelocityValues(real* vx, real* vy, real* vz, int* indi { indices[allIndicesCounter] = grids[level]->getSparseIndex(boundaryCondition->indices[i]) +1; - vx[allIndicesCounter] = boundaryCondition->getVx(i); - vy[allIndicesCounter] = boundaryCondition->getVy(i); - vz[allIndicesCounter] = boundaryCondition->getVz(i); + vx[allIndicesCounter] = (uint)boundaryCondition->getVx(i); + vy[allIndicesCounter] = (uint)boundaryCondition->getVy(i); + vz[allIndicesCounter] = (uint)boundaryCondition->getVz(i); allIndicesCounter++; } } @@ -419,7 +419,7 @@ void LevelGridBuilder::getPressureQs(real* qs[27], int level) const uint LevelGridBuilder::getGeometrySize(int level) const { if (boundaryConditions[level]->geometryBoundaryCondition) - return boundaryConditions[level]->geometryBoundaryCondition->indices.size(); + return (uint)boundaryConditions[level]->geometryBoundaryCondition->indices.size(); return 0; } diff --git a/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp b/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp index 86d3c2b71..8c36de45c 100644 --- a/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp +++ b/src/gpu/GridGenerator/grid/GridBuilder/MultipleGridBuilder.cpp @@ -461,9 +461,9 @@ void MultipleGridBuilder::buildGrids( LbmOrGks lbmOrGks, bool enableThinWalls ) // Figure 5.2 in the Dissertation of Stephan Lenz: // https://publikationsserver.tu-braunschweig.de/receive/dbbs_mods_00068716 // - for( int level = grids.size()-1; level >= 0; level-- ) { + for( std::size_t level = grids.size()-1; level >= 0; level-- ) { - *logging::out << logging::Logger::INFO_INTERMEDIATE << "Start initializing level " << level << "\n"; + *logging::out << logging::Logger::INFO_INTERMEDIATE << "Start initializing level " << (int)level << "\n"; // On the coarse grid every thing is Fluid (w.r.t. the refinement) // On the finest grid the Fluid region is defined by the Object @@ -475,7 +475,7 @@ void MultipleGridBuilder::buildGrids( LbmOrGks lbmOrGks, bool enableThinWalls ) else grids[level]->inital( grids[level+1], this->numberOfLayersBetweenLevels ); - *logging::out << logging::Logger::INFO_INTERMEDIATE << "Done initializing level " << level << "\n"; + *logging::out << logging::Logger::INFO_INTERMEDIATE << "Done initializing level " << (int)level << "\n"; } ////////////////////////////////////////////////////////////////////////// @@ -500,7 +500,7 @@ void MultipleGridBuilder::buildGrids( LbmOrGks lbmOrGks, bool enableThinWalls ) // change the following two lines. This is not tested though! //for( uint level = 0; level < grids.size(); level++ ) - uint level = grids.size() - 1; + uint level = (uint)grids.size() - 1; { // the Grid::mesh(...) method distinguishes inside and ouside regions // of the solid domain.: diff --git a/src/gpu/GridGenerator/grid/GridImp.cu b/src/gpu/GridGenerator/grid/GridImp.cu index 61bf1c75d..cd835ee1c 100644 --- a/src/gpu/GridGenerator/grid/GridImp.cu +++ b/src/gpu/GridGenerator/grid/GridImp.cu @@ -180,7 +180,7 @@ HOSTDEVICE void GridImp::findInnerNode(uint index) HOSTDEVICE void GridImp::discretize(Object* solidObject, char innerType, char outerType) { #pragma omp parallel for - for (int index = 0; index < this->size; index++) + for (int index = 0; index < (int)this->size; index++) { this->sparseIndices[index] = index; @@ -1299,7 +1299,7 @@ CUDA_HOST void GridImp::findQsPrimitive(Object * object) } - for( int index = 0; index < this->size; index++ ) + for( int index = 0; index < (int)this->size; index++ ) { if( this->qIndices[index] == INVALID_INDEX ) continue; @@ -1554,12 +1554,12 @@ void GridImp::findCommunicationIndex( uint index, real coordinate, real limit, i uint GridImp::getNumberOfSendNodes(int direction) { - return this->communicationIndices[direction].sendIndices.size(); + return (uint)this->communicationIndices[direction].sendIndices.size(); } uint GridImp::getNumberOfReceiveNodes(int direction) { - return this->communicationIndices[direction].receiveIndices.size(); + return (uint)this->communicationIndices[direction].receiveIndices.size(); } uint GridImp::getSendIndex(int direction, uint index) diff --git a/src/gpu/GridGenerator/grid/GridImp.h b/src/gpu/GridGenerator/grid/GridImp.h index 6cb640682..314020d0e 100644 --- a/src/gpu/GridGenerator/grid/GridImp.h +++ b/src/gpu/GridGenerator/grid/GridImp.h @@ -28,7 +28,7 @@ class TriangularMeshDiscretizationStrategy; #endif #endif -// warning #3156-D: extern declaration of the entity DIRECTIONS is treated as a static definition +//GCC: warning #3156-D: extern declaration of the entity DIRECTIONS is treated as a static definition extern CONSTANT int DIRECTIONS[DIR_END_MAX][DIMENSION]; #ifdef __GNUC__ diff --git a/src/gpu/GridGenerator/grid/GridInterface.cu b/src/gpu/GridGenerator/grid/GridInterface.cu index 00ce04f09..ec88c3b4a 100644 --- a/src/gpu/GridGenerator/grid/GridInterface.cu +++ b/src/gpu/GridGenerator/grid/GridInterface.cu @@ -278,7 +278,7 @@ CUDA_HOST void GRIDGENERATOR_EXPORT GridInterface::repairGridInterfaceOnMultiGPU delete[] cf.fine; delete[] cf.offset; - cf.numberOfEntries = tmpCFC.size(); + cf.numberOfEntries = (uint)tmpCFC.size(); cf.coarse = new uint[cf.numberOfEntries]; cf.fine = new uint[cf.numberOfEntries]; @@ -311,7 +311,7 @@ CUDA_HOST void GRIDGENERATOR_EXPORT GridInterface::repairGridInterfaceOnMultiGPU delete[] fc.coarse; delete[] fc.offset; - fc.numberOfEntries = tmpFCC.size(); + fc.numberOfEntries = (uint)tmpFCC.size(); fc.fine = new uint[fc.numberOfEntries]; fc.coarse = new uint[fc.numberOfEntries]; diff --git a/src/gpu/GridGenerator/grid/GridStrategy/GridCpuStrategy/GridCpuStrategy.cpp b/src/gpu/GridGenerator/grid/GridStrategy/GridCpuStrategy/GridCpuStrategy.cpp index c906ad36b..34fbe03eb 100644 --- a/src/gpu/GridGenerator/grid/GridStrategy/GridCpuStrategy/GridCpuStrategy.cpp +++ b/src/gpu/GridGenerator/grid/GridStrategy/GridCpuStrategy/GridCpuStrategy.cpp @@ -208,18 +208,18 @@ void GridCpuStrategy::findGridInterface(SPtr<GridImp> grid, SPtr<GridImp> fineGr grid->gridInterface->fc.fine = new uint[sizeCF]; grid->gridInterface->fc.offset = new uint[sizeCF]; - for (std::size_t index = 0; index < grid->getSize(); index++) + for (uint index = 0; index < grid->getSize(); index++) grid->findGridInterfaceCF(index, *fineGrid, lbmOrGks); - for (std::size_t index = 0; index < grid->getSize(); index++) + for (uint index = 0; index < grid->getSize(); index++) grid->findGridInterfaceFC(index, *fineGrid); - for (std::size_t index = 0; index < grid->getSize(); index++) + for (uint index = 0; index < grid->getSize(); index++) grid->findOverlapStopper(index, *fineGrid); if( lbmOrGks == GKS ) { - for (std::size_t index = 0; index < grid->getSize(); index++) + for (uint index = 0; index < grid->getSize(); index++) grid->findInvalidBoundaryNodes(index); } diff --git a/src/gpu/GridGenerator/io/STLReaderWriter/STLReader.cpp b/src/gpu/GridGenerator/io/STLReaderWriter/STLReader.cpp index 69cd83961..3f41f66ae 100644 --- a/src/gpu/GridGenerator/io/STLReaderWriter/STLReader.cpp +++ b/src/gpu/GridGenerator/io/STLReaderWriter/STLReader.cpp @@ -295,7 +295,7 @@ std::vector<Triangle> STLReader::readBinarySTL(const BoundingBox &box, const std if (box.isInside(t) || box.intersect(t)) triangles.push_back(t); } - int size = triangles.size(); + int size = (int)triangles.size(); *logging::out << logging::Logger::INFO_INTERMEDIATE << "Number of Triangles in process: " << size << "\n"; *logging::out << logging::Logger::INFO_INTERMEDIATE << "Complete reading STL file. \n"; diff --git a/src/gpu/GridGenerator/io/STLReaderWriter/STLWriter.cpp b/src/gpu/GridGenerator/io/STLReaderWriter/STLWriter.cpp index 724ae89b0..36e099c61 100644 --- a/src/gpu/GridGenerator/io/STLReaderWriter/STLWriter.cpp +++ b/src/gpu/GridGenerator/io/STLReaderWriter/STLWriter.cpp @@ -9,7 +9,7 @@ void STLWriter::writeSTL(std::vector<Triangle> &vec, const std::string &name, bool writeBinary) { - const int size = vec.size(); + const int size = (int)vec.size(); *logging::out << logging::Logger::INFO_INTERMEDIATE << "Write " << size << " Triangles to STL : " + name + "\n"; std::ios_base::openmode mode = std::ios::out; diff --git a/src/gpu/GridGenerator/io/SimulationFileWriter/SimulationFileWriter.cpp b/src/gpu/GridGenerator/io/SimulationFileWriter/SimulationFileWriter.cpp index ed1403dd8..f3a66f69a 100644 --- a/src/gpu/GridGenerator/io/SimulationFileWriter/SimulationFileWriter.cpp +++ b/src/gpu/GridGenerator/io/SimulationFileWriter/SimulationFileWriter.cpp @@ -575,7 +575,7 @@ void SimulationFileWriter::writeBoundaryShort(std::vector<real> boundary, int rb void SimulationFileWriter::writeBoundaryShort(SPtr<Grid> grid, SPtr<BoundaryCondition> boundaryCondition, uint side) { - uint numberOfBoundaryNodes = boundaryCondition->indices.size(); + uint numberOfBoundaryNodes = (uint)boundaryCondition->indices.size(); *valueStreams[side] << numberOfBoundaryNodes << "\n"; *qStreams[side] << numberOfBoundaryNodes << "\n"; diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp index 556540a16..49de1d53a 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp @@ -379,13 +379,13 @@ void GridReader::initalValuesDomainDecompostion(int level) para->getParD(i)->recvProcessNeighborX[j].memsizeFs = sizeof(real) *tempRecv; //////////////////////////////////////////////////////////////////////////////////////// //malloc on host and device - cudaMemoryManager->cudaAllocProcessNeighborX(i, j); + cudaMemoryManager->cudaAllocProcessNeighborX(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// //init index arrays procNeighborsSendX[j]->initIndex(para->getParH(i)->sendProcessNeighborX[j].index, i); procNeighborsRecvX[j]->initIndex(para->getParH(i)->recvProcessNeighborX[j].index, i); //////////////////////////////////////////////////////////////////////////////////////// - cudaMemoryManager->cudaCopyProcessNeighborXIndex(i, j); + cudaMemoryManager->cudaCopyProcessNeighborXIndex(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// } } @@ -435,13 +435,13 @@ void GridReader::initalValuesDomainDecompostion(int level) para->getParD(i)->recvProcessNeighborY[j].memsizeFs = sizeof(real) *tempRecv; //////////////////////////////////////////////////////////////////////////////////////// //malloc on host and device - cudaMemoryManager->cudaAllocProcessNeighborY(i, j); + cudaMemoryManager->cudaAllocProcessNeighborY(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// //init index arrays procNeighborsSendY[j]->initIndex(para->getParH(i)->sendProcessNeighborY[j].index, i); procNeighborsRecvY[j]->initIndex(para->getParH(i)->recvProcessNeighborY[j].index, i); //////////////////////////////////////////////////////////////////////////////////////// - cudaMemoryManager->cudaCopyProcessNeighborYIndex(i, j); + cudaMemoryManager->cudaCopyProcessNeighborYIndex(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// } } @@ -491,13 +491,13 @@ void GridReader::initalValuesDomainDecompostion(int level) para->getParD(i)->recvProcessNeighborZ[j].memsizeFs = sizeof(real) *tempRecv; //////////////////////////////////////////////////////////////////////////////////////// //malloc on host and device - cudaMemoryManager->cudaAllocProcessNeighborZ(i, j); + cudaMemoryManager->cudaAllocProcessNeighborZ(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// //init index arrays procNeighborsSendZ[j]->initIndex(para->getParH(i)->sendProcessNeighborZ[j].index, i); procNeighborsRecvZ[j]->initIndex(para->getParH(i)->recvProcessNeighborZ[j].index, i); //////////////////////////////////////////////////////////////////////////////////////// - cudaMemoryManager->cudaCopyProcessNeighborZIndex(i, j); + cudaMemoryManager->cudaCopyProcessNeighborZIndex(i, (uint)j); //////////////////////////////////////////////////////////////////////////////////////// } } diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp index 4c20ea485..2582e6918 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp @@ -253,7 +253,7 @@ void GridGenerator::allocArrays_BoundaryValues() { if( direction == CommunicationDirections::MX || direction == CommunicationDirections::PX ) { - int j = para->getParH(level)->sendProcessNeighborX.size(); + int j = (int)para->getParH(level)->sendProcessNeighborX.size(); para->getParH(level)->sendProcessNeighborX.emplace_back(); para->getParD(level)->sendProcessNeighborX.emplace_back(); @@ -313,7 +313,7 @@ void GridGenerator::allocArrays_BoundaryValues() if( direction == CommunicationDirections::MY || direction == CommunicationDirections::PY ) { - int j = para->getParH(level)->sendProcessNeighborY.size(); + int j = (int)para->getParH(level)->sendProcessNeighborY.size(); para->getParH(level)->sendProcessNeighborY.emplace_back(); para->getParD(level)->sendProcessNeighborY.emplace_back(); @@ -373,7 +373,7 @@ void GridGenerator::allocArrays_BoundaryValues() if( direction == CommunicationDirections::MZ || direction == CommunicationDirections::PZ ) { - int j = para->getParH(level)->sendProcessNeighborZ.size(); + int j = (int)para->getParH(level)->sendProcessNeighborZ.size(); para->getParH(level)->sendProcessNeighborZ.emplace_back(); para->getParD(level)->sendProcessNeighborZ.emplace_back(); @@ -447,7 +447,7 @@ void GridGenerator::allocArrays_BoundaryValues() { if (direction == CommunicationDirections::MX || direction == CommunicationDirections::PX) { - int j = para->getParH(level)->sendProcessNeighborF3X.size(); + int j = (int)para->getParH(level)->sendProcessNeighborF3X.size(); para->getParH(level)->sendProcessNeighborF3X.emplace_back(); para->getParD(level)->sendProcessNeighborF3X.emplace_back(); @@ -501,7 +501,7 @@ void GridGenerator::allocArrays_BoundaryValues() if (direction == CommunicationDirections::MY || direction == CommunicationDirections::PY) { - int j = para->getParH(level)->sendProcessNeighborF3Y.size(); + int j = (int)para->getParH(level)->sendProcessNeighborF3Y.size(); para->getParH(level)->sendProcessNeighborF3Y.emplace_back(); para->getParD(level)->sendProcessNeighborF3Y.emplace_back(); @@ -555,7 +555,7 @@ void GridGenerator::allocArrays_BoundaryValues() if (direction == CommunicationDirections::MZ || direction == CommunicationDirections::PZ) { - int j = para->getParH(level)->sendProcessNeighborF3Z.size(); + int j = (int)para->getParH(level)->sendProcessNeighborF3Z.size(); para->getParH(level)->sendProcessNeighborF3Z.emplace_back(); para->getParD(level)->sendProcessNeighborF3Z.emplace_back(); @@ -820,9 +820,9 @@ void GridGenerator::allocArrays_BoundaryQs() builder->getGeometryQs(Q.q27, i); //QDebugWriter::writeQValues(Q, para->getParH(i)->QGeom.k, para->getParH(i)->QGeom.kQ, "M:/TestGridGeneration/results/GeomGPU.dat"); ////////////////////////////////////////////////////////////////// - for (int i = 0; i < numberOfGeometryNodes; i++) + for (int node_i = 0; node_i < numberOfGeometryNodes; node_i++) { - Q.q27[dirZERO][i] = 0.0f; + Q.q27[dirZERO][node_i] = 0.0f; } //for(int test = 0; test < 3; test++) //{ diff --git a/src/gpu/VirtualFluids_GPU/GPU/AdvecDiffBCs27.cu b/src/gpu/VirtualFluids_GPU/GPU/AdvecDiffBCs27.cu index b044bb70f..6e67ecda7 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/AdvecDiffBCs27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/AdvecDiffBCs27.cu @@ -2130,67 +2130,67 @@ extern "C" __global__ void QADVel27(int inx, vx2 = OORho*((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); vx3 = OORho*((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //////////////////////////////////////////////////////////////////////////////// - real f27_W = (D27.f[dirE ])[ke ]; - real f27_E = (D27.f[dirW ])[kw ]; - real f27_S = (D27.f[dirN ])[kn ]; - real f27_N = (D27.f[dirS ])[ks ]; - real f27_B = (D27.f[dirT ])[kt ]; - real f27_T = (D27.f[dirB ])[kb ]; - real f27_SW = (D27.f[dirNE ])[kne ]; - real f27_NE = (D27.f[dirSW ])[ksw ]; - real f27_NW = (D27.f[dirSE ])[kse ]; - real f27_SE = (D27.f[dirNW ])[knw ]; - real f27_BW = (D27.f[dirTE ])[kte ]; - real f27_TE = (D27.f[dirBW ])[kbw ]; - real f27_TW = (D27.f[dirBE ])[kbe ]; - real f27_BE = (D27.f[dirTW ])[ktw ]; - real f27_BS = (D27.f[dirTN ])[ktn ]; - real f27_TN = (D27.f[dirBS ])[kbs ]; - real f27_TS = (D27.f[dirBN ])[kbn ]; - real f27_BN = (D27.f[dirTS ])[kts ]; - real f27_ZERO = (D27.f[dirZERO])[kzero]; - real f27_BSW = (D27.f[dirTNE ])[ktne ]; - real f27_BNE = (D27.f[dirTSW ])[ktsw ]; - real f27_BNW = (D27.f[dirTSE ])[ktse ]; - real f27_BSE = (D27.f[dirTNW ])[ktnw ]; - real f27_TSW = (D27.f[dirBNE ])[kbne ]; - real f27_TNE = (D27.f[dirBSW ])[kbsw ]; - real f27_TNW = (D27.f[dirBSE ])[kbse ]; - real f27_TSE = (D27.f[dirBNW ])[kbnw ]; + //real f27_W = (D27.f[dirE ])[ke ]; + //real f27_E = (D27.f[dirW ])[kw ]; + //real f27_S = (D27.f[dirN ])[kn ]; + //real f27_N = (D27.f[dirS ])[ks ]; + //real f27_B = (D27.f[dirT ])[kt ]; + //real f27_T = (D27.f[dirB ])[kb ]; + //real f27_SW = (D27.f[dirNE ])[kne ]; + //real f27_NE = (D27.f[dirSW ])[ksw ]; + //real f27_NW = (D27.f[dirSE ])[kse ]; + //real f27_SE = (D27.f[dirNW ])[knw ]; + //real f27_BW = (D27.f[dirTE ])[kte ]; + //real f27_TE = (D27.f[dirBW ])[kbw ]; + //real f27_TW = (D27.f[dirBE ])[kbe ]; + //real f27_BE = (D27.f[dirTW ])[ktw ]; + //real f27_BS = (D27.f[dirTN ])[ktn ]; + //real f27_TN = (D27.f[dirBS ])[kbs ]; + //real f27_TS = (D27.f[dirBN ])[kbn ]; + //real f27_BN = (D27.f[dirTS ])[kts ]; + //real f27_ZERO = (D27.f[dirZERO])[kzero]; + //real f27_BSW = (D27.f[dirTNE ])[ktne ]; + //real f27_BNE = (D27.f[dirTSW ])[ktsw ]; + //real f27_BNW = (D27.f[dirTSE ])[ktse ]; + //real f27_BSE = (D27.f[dirTNW ])[ktnw ]; + //real f27_TSW = (D27.f[dirBNE ])[kbne ]; + //real f27_TNE = (D27.f[dirBSW ])[kbsw ]; + //real f27_TNW = (D27.f[dirBSE ])[kbse ]; + //real f27_TSE = (D27.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); //////////////////////////////////////////////////////////////////////////////// - real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + - f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + - f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; + //real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + + // f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + + // f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; //real feq27_ZERO = c8over27* ConcD*(one-cu_sq); - real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); - real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); - real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); - real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); - real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); - real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); - real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); - real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); - real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); - real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); - real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); - real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); - real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); - real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); - real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); - real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); - real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); - real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); - real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); - real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); - real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); - real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); - real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); - real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); - real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); + //real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); + //real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); + //real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); + //real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); + //real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); + //real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); + //real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); + //real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); + //real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); + //real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); + //real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); + //real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); + //real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); + //real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); + //real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); + //real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); + //real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); + //real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); + //real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); + //real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); + //real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); + //real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); + //real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); + //real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); + //real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); + //real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// real TempD = temp[k]; //real TempD = four; @@ -3432,67 +3432,67 @@ extern "C" __global__ void QADBB27(int inx, unsigned int size_Mat, bool evenOrOdd) { - Distributions27 D; - if (evenOrOdd==true) - { - D.f[dirE ] = &DD[dirE *size_Mat]; - D.f[dirW ] = &DD[dirW *size_Mat]; - D.f[dirN ] = &DD[dirN *size_Mat]; - D.f[dirS ] = &DD[dirS *size_Mat]; - D.f[dirT ] = &DD[dirT *size_Mat]; - D.f[dirB ] = &DD[dirB *size_Mat]; - D.f[dirNE ] = &DD[dirNE *size_Mat]; - D.f[dirSW ] = &DD[dirSW *size_Mat]; - D.f[dirSE ] = &DD[dirSE *size_Mat]; - D.f[dirNW ] = &DD[dirNW *size_Mat]; - D.f[dirTE ] = &DD[dirTE *size_Mat]; - D.f[dirBW ] = &DD[dirBW *size_Mat]; - D.f[dirBE ] = &DD[dirBE *size_Mat]; - D.f[dirTW ] = &DD[dirTW *size_Mat]; - D.f[dirTN ] = &DD[dirTN *size_Mat]; - D.f[dirBS ] = &DD[dirBS *size_Mat]; - D.f[dirBN ] = &DD[dirBN *size_Mat]; - D.f[dirTS ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirTNE *size_Mat]; - D.f[dirTSW ] = &DD[dirTSW *size_Mat]; - D.f[dirTSE ] = &DD[dirTSE *size_Mat]; - D.f[dirTNW ] = &DD[dirTNW *size_Mat]; - D.f[dirBNE ] = &DD[dirBNE *size_Mat]; - D.f[dirBSW ] = &DD[dirBSW *size_Mat]; - D.f[dirBSE ] = &DD[dirBSE *size_Mat]; - D.f[dirBNW ] = &DD[dirBNW *size_Mat]; - } - else - { - D.f[dirW ] = &DD[dirE *size_Mat]; - D.f[dirE ] = &DD[dirW *size_Mat]; - D.f[dirS ] = &DD[dirN *size_Mat]; - D.f[dirN ] = &DD[dirS *size_Mat]; - D.f[dirB ] = &DD[dirT *size_Mat]; - D.f[dirT ] = &DD[dirB *size_Mat]; - D.f[dirSW ] = &DD[dirNE *size_Mat]; - D.f[dirNE ] = &DD[dirSW *size_Mat]; - D.f[dirNW ] = &DD[dirSE *size_Mat]; - D.f[dirSE ] = &DD[dirNW *size_Mat]; - D.f[dirBW ] = &DD[dirTE *size_Mat]; - D.f[dirTE ] = &DD[dirBW *size_Mat]; - D.f[dirTW ] = &DD[dirBE *size_Mat]; - D.f[dirBE ] = &DD[dirTW *size_Mat]; - D.f[dirBS ] = &DD[dirTN *size_Mat]; - D.f[dirTN ] = &DD[dirBS *size_Mat]; - D.f[dirTS ] = &DD[dirBN *size_Mat]; - D.f[dirBN ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirBSW *size_Mat]; - D.f[dirTSW ] = &DD[dirBNE *size_Mat]; - D.f[dirTSE ] = &DD[dirBNW *size_Mat]; - D.f[dirTNW ] = &DD[dirBSE *size_Mat]; - D.f[dirBNE ] = &DD[dirTSW *size_Mat]; - D.f[dirBSW ] = &DD[dirTNE *size_Mat]; - D.f[dirBSE ] = &DD[dirTNW *size_Mat]; - D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } + //Distributions27 D; + //if (evenOrOdd==true) + //{ + // D.f[dirE ] = &DD[dirE *size_Mat]; + // D.f[dirW ] = &DD[dirW *size_Mat]; + // D.f[dirN ] = &DD[dirN *size_Mat]; + // D.f[dirS ] = &DD[dirS *size_Mat]; + // D.f[dirT ] = &DD[dirT *size_Mat]; + // D.f[dirB ] = &DD[dirB *size_Mat]; + // D.f[dirNE ] = &DD[dirNE *size_Mat]; + // D.f[dirSW ] = &DD[dirSW *size_Mat]; + // D.f[dirSE ] = &DD[dirSE *size_Mat]; + // D.f[dirNW ] = &DD[dirNW *size_Mat]; + // D.f[dirTE ] = &DD[dirTE *size_Mat]; + // D.f[dirBW ] = &DD[dirBW *size_Mat]; + // D.f[dirBE ] = &DD[dirBE *size_Mat]; + // D.f[dirTW ] = &DD[dirTW *size_Mat]; + // D.f[dirTN ] = &DD[dirTN *size_Mat]; + // D.f[dirBS ] = &DD[dirBS *size_Mat]; + // D.f[dirBN ] = &DD[dirBN *size_Mat]; + // D.f[dirTS ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirTNE *size_Mat]; + // D.f[dirTSW ] = &DD[dirTSW *size_Mat]; + // D.f[dirTSE ] = &DD[dirTSE *size_Mat]; + // D.f[dirTNW ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNE ] = &DD[dirBNE *size_Mat]; + // D.f[dirBSW ] = &DD[dirBSW *size_Mat]; + // D.f[dirBSE ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNW ] = &DD[dirBNW *size_Mat]; + //} + //else + //{ + // D.f[dirW ] = &DD[dirE *size_Mat]; + // D.f[dirE ] = &DD[dirW *size_Mat]; + // D.f[dirS ] = &DD[dirN *size_Mat]; + // D.f[dirN ] = &DD[dirS *size_Mat]; + // D.f[dirB ] = &DD[dirT *size_Mat]; + // D.f[dirT ] = &DD[dirB *size_Mat]; + // D.f[dirSW ] = &DD[dirNE *size_Mat]; + // D.f[dirNE ] = &DD[dirSW *size_Mat]; + // D.f[dirNW ] = &DD[dirSE *size_Mat]; + // D.f[dirSE ] = &DD[dirNW *size_Mat]; + // D.f[dirBW ] = &DD[dirTE *size_Mat]; + // D.f[dirTE ] = &DD[dirBW *size_Mat]; + // D.f[dirTW ] = &DD[dirBE *size_Mat]; + // D.f[dirBE ] = &DD[dirTW *size_Mat]; + // D.f[dirBS ] = &DD[dirTN *size_Mat]; + // D.f[dirTN ] = &DD[dirBS *size_Mat]; + // D.f[dirTS ] = &DD[dirBN *size_Mat]; + // D.f[dirBN ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirBSW *size_Mat]; + // D.f[dirTSW ] = &DD[dirBNE *size_Mat]; + // D.f[dirTSE ] = &DD[dirBNW *size_Mat]; + // D.f[dirTNW ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNE ] = &DD[dirTSW *size_Mat]; + // D.f[dirBSW ] = &DD[dirTNE *size_Mat]; + // D.f[dirBSE ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNW ] = &DD[dirTSE *size_Mat]; + //} Distributions27 D27; if (evenOrOdd==true) @@ -3603,7 +3603,7 @@ extern "C" __global__ void QADBB27(int inx, //////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; + //unsigned int kzero= KQK; unsigned int ke = KQK; unsigned int kw = neighborX[KQK]; unsigned int kn = KQK; @@ -3631,35 +3631,36 @@ extern "C" __global__ void QADBB27(int inx, unsigned int ktne = KQK; unsigned int kbsw = neighborZ[ksw]; //////////////////////////////////////////////////////////////////////////////// - real f_W = (D.f[dirE ])[ke ]; - real f_E = (D.f[dirW ])[kw ]; - real f_S = (D.f[dirN ])[kn ]; - real f_N = (D.f[dirS ])[ks ]; - real f_B = (D.f[dirT ])[kt ]; - real f_T = (D.f[dirB ])[kb ]; - real f_SW = (D.f[dirNE ])[kne ]; - real f_NE = (D.f[dirSW ])[ksw ]; - real f_NW = (D.f[dirSE ])[kse ]; - real f_SE = (D.f[dirNW ])[knw ]; - real f_BW = (D.f[dirTE ])[kte ]; - real f_TE = (D.f[dirBW ])[kbw ]; - real f_TW = (D.f[dirBE ])[kbe ]; - real f_BE = (D.f[dirTW ])[ktw ]; - real f_BS = (D.f[dirTN ])[ktn ]; - real f_TN = (D.f[dirBS ])[kbs ]; - real f_TS = (D.f[dirBN ])[kbn ]; - real f_BN = (D.f[dirTS ])[kts ]; - real f_ZERO = (D.f[dirZERO])[kzero]; - real f_BSW = (D.f[dirTNE ])[ktne ]; - real f_BNE = (D.f[dirTSW ])[ktsw ]; - real f_BNW = (D.f[dirTSE ])[ktse ]; - real f_BSE = (D.f[dirTNW ])[ktnw ]; - real f_TSW = (D.f[dirBNE ])[kbne ]; - real f_TNE = (D.f[dirBSW ])[kbsw ]; - real f_TNW = (D.f[dirBSE ])[kbse ]; - real f_TSE = (D.f[dirBNW ])[kbnw ]; + //real f_W = (D.f[dirE ])[ke ]; + //real f_E = (D.f[dirW ])[kw ]; + //real f_S = (D.f[dirN ])[kn ]; + //real f_N = (D.f[dirS ])[ks ]; + //real f_B = (D.f[dirT ])[kt ]; + //real f_T = (D.f[dirB ])[kb ]; + //real f_SW = (D.f[dirNE ])[kne ]; + //real f_NE = (D.f[dirSW ])[ksw ]; + //real f_NW = (D.f[dirSE ])[kse ]; + //real f_SE = (D.f[dirNW ])[knw ]; + //real f_BW = (D.f[dirTE ])[kte ]; + //real f_TE = (D.f[dirBW ])[kbw ]; + //real f_TW = (D.f[dirBE ])[kbe ]; + //real f_BE = (D.f[dirTW ])[ktw ]; + //real f_BS = (D.f[dirTN ])[ktn ]; + //real f_TN = (D.f[dirBS ])[kbs ]; + //real f_TS = (D.f[dirBN ])[kbn ]; + //real f_BN = (D.f[dirTS ])[kts ]; + //real f_ZERO = (D.f[dirZERO])[kzero]; + //real f_BSW = (D.f[dirTNE ])[ktne ]; + //real f_BNE = (D.f[dirTSW ])[ktsw ]; + //real f_BNW = (D.f[dirTSE ])[ktse ]; + //real f_BSE = (D.f[dirTNW ])[ktnw ]; + //real f_TSW = (D.f[dirBNE ])[kbne ]; + //real f_TNE = (D.f[dirBSW ])[kbsw ]; + //real f_TNW = (D.f[dirBSE ])[kbse ]; + //real f_TSE = (D.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// - real vx1, vx2, vx3, /*drho, feq,*/ q; + //real vx1, vx2, vx3, /*drho, feq,*/ q; + real q; ////drho = f_TSE + f_TNW + f_TNE + f_TSW + f_BSE + f_BNW + f_BNE + f_BSW + //// f_BN + f_TS + f_TN + f_BS + f_BE + f_TW + f_TE + f_BW + f_SE + f_NW + f_NE + f_SW + //// f_T + f_B + f_N + f_S + f_E + f_W + f_ZERO; @@ -3676,12 +3677,12 @@ extern "C" __global__ void QADBB27(int inx, //vx3 = ((f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) + (f_TSW - f_BNE)) + // (-(f_BN - f_TS) + (f_TN - f_BS)) + ((f_TE - f_BW) - (f_BE - f_TW)) + // (f_T - f_B); - real rho0 = (f_TNE+f_BSW)+(f_TSW+f_BNE)+(f_TSE+f_BNW)+(f_TNW+f_BSE)+(f_NE+f_SW)+(f_NW+f_SE)+(f_TE+f_BW)+(f_BE+f_TW)+(f_TN+f_BS)+(f_BN+f_TS)+(f_E+f_W)+(f_N+f_S)+(f_T+f_B)+f_ZERO; - real rho = rho0 + c1o1; - real OORho = c1o1/rho; - vx1 = OORho*((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); - vx2 = OORho*((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); - vx3 = OORho*((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); + //real rho0 = (f_TNE+f_BSW)+(f_TSW+f_BNE)+(f_TSE+f_BNW)+(f_TNW+f_BSE)+(f_NE+f_SW)+(f_NW+f_SE)+(f_TE+f_BW)+(f_BE+f_TW)+(f_TN+f_BS)+(f_BN+f_TS)+(f_E+f_W)+(f_N+f_S)+(f_T+f_B)+f_ZERO; + //real rho = rho0 + c1o1; + //real OORho = c1o1/rho; + //vx1 = OORho*((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); + //vx2 = OORho*((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); + //vx3 = OORho*((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //////////////////////////////////////////////////////////////////////////////// real f27_W = (D27.f[dirE ])[ke ]; real f27_E = (D27.f[dirW ])[kw ]; @@ -3701,7 +3702,7 @@ extern "C" __global__ void QADBB27(int inx, real f27_TN = (D27.f[dirBS ])[kbs ]; real f27_TS = (D27.f[dirBN ])[kbn ]; real f27_BN = (D27.f[dirTS ])[kts ]; - real f27_ZERO = (D27.f[dirZERO])[kzero]; + //real f27_ZERO = (D27.f[dirZERO])[kzero]; real f27_BSW = (D27.f[dirTNE ])[ktne ]; real f27_BNE = (D27.f[dirTSW ])[ktsw ]; real f27_BNW = (D27.f[dirTSE ])[ktse ]; @@ -3711,69 +3712,69 @@ extern "C" __global__ void QADBB27(int inx, real f27_TNW = (D27.f[dirBSE ])[kbse ]; real f27_TSE = (D27.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// - real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); + //real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); //////////////////////////////////////////////////////////////////////////////// - real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + - f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + - f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; + //real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + + // f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + + // f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; //real feq27_ZERO = c8over27* ConcD*(one-cu_sq); - real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); - real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); - real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); - real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); - real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); - real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); - real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); - real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); - real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); - real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); - real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); - real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); - real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); - real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); - real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); - real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); - real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); - real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); - real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); - real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); - real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); - real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); - real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); - real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); - real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); + //real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); + //real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); + //real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); + //real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); + //real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); + //real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); + //real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); + //real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); + //real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); + //real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); + //real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); + //real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); + //real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); + //real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); + //real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); + //real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); + //real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); + //real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); + //real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); + //real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); + //real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); + //real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); + //real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); + //real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); + //real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); + //real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - real TempD = temp[k]; + //real TempD = temp[k]; //real feqW27_ZERO = c8over27* TempD*(one-cu_sq); - real feqW27_E = c2o27* TempD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); - real feqW27_W = c2o27* TempD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); - real feqW27_N = c2o27* TempD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); - real feqW27_S = c2o27* TempD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); - real feqW27_T = c2o27* TempD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); - real feqW27_B = c2o27* TempD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); - real feqW27_NE = c1o54* TempD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); - real feqW27_SW = c1o54* TempD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); - real feqW27_SE = c1o54* TempD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); - real feqW27_NW = c1o54* TempD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); - real feqW27_TE = c1o54* TempD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); - real feqW27_BW = c1o54* TempD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); - real feqW27_BE = c1o54* TempD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); - real feqW27_TW = c1o54* TempD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); - real feqW27_TN = c1o54* TempD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); - real feqW27_BS = c1o54* TempD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); - real feqW27_BN = c1o54* TempD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); - real feqW27_TS = c1o54* TempD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); - real feqW27_TNE = c1o216*TempD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); - real feqW27_BSW = c1o216*TempD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); - real feqW27_BNE = c1o216*TempD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); - real feqW27_TSW = c1o216*TempD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); - real feqW27_TSE = c1o216*TempD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); - real feqW27_BNW = c1o216*TempD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); - real feqW27_BSE = c1o216*TempD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - real feqW27_TNW = c1o216*TempD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); + //real feqW27_E = c2o27* TempD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); + //real feqW27_W = c2o27* TempD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); + //real feqW27_N = c2o27* TempD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); + //real feqW27_S = c2o27* TempD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); + //real feqW27_T = c2o27* TempD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); + //real feqW27_B = c2o27* TempD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); + //real feqW27_NE = c1o54* TempD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); + //real feqW27_SW = c1o54* TempD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); + //real feqW27_SE = c1o54* TempD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); + //real feqW27_NW = c1o54* TempD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); + //real feqW27_TE = c1o54* TempD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); + //real feqW27_BW = c1o54* TempD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); + //real feqW27_BE = c1o54* TempD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); + //real feqW27_TW = c1o54* TempD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); + //real feqW27_TN = c1o54* TempD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); + //real feqW27_BS = c1o54* TempD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); + //real feqW27_BN = c1o54* TempD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); + //real feqW27_TS = c1o54* TempD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); + //real feqW27_TNE = c1o216*TempD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); + //real feqW27_BSW = c1o216*TempD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); + //real feqW27_BNE = c1o216*TempD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); + //real feqW27_TSW = c1o216*TempD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); + //real feqW27_TSE = c1o216*TempD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); + //real feqW27_BNW = c1o216*TempD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); + //real feqW27_BSE = c1o216*TempD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); + //real feqW27_TNW = c1o216*TempD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// real omegaD = c3o1 - sqrt(c3o1); //real Lam = -(c1o2-one/omegaD); @@ -3949,67 +3950,67 @@ extern "C" __global__ void QNoSlipADincomp7( int inx, unsigned int size_Mat, bool evenOrOdd) { - Distributions27 D; - if (evenOrOdd==true) - { - D.f[dirE ] = &DD[dirE *size_Mat]; - D.f[dirW ] = &DD[dirW *size_Mat]; - D.f[dirN ] = &DD[dirN *size_Mat]; - D.f[dirS ] = &DD[dirS *size_Mat]; - D.f[dirT ] = &DD[dirT *size_Mat]; - D.f[dirB ] = &DD[dirB *size_Mat]; - D.f[dirNE ] = &DD[dirNE *size_Mat]; - D.f[dirSW ] = &DD[dirSW *size_Mat]; - D.f[dirSE ] = &DD[dirSE *size_Mat]; - D.f[dirNW ] = &DD[dirNW *size_Mat]; - D.f[dirTE ] = &DD[dirTE *size_Mat]; - D.f[dirBW ] = &DD[dirBW *size_Mat]; - D.f[dirBE ] = &DD[dirBE *size_Mat]; - D.f[dirTW ] = &DD[dirTW *size_Mat]; - D.f[dirTN ] = &DD[dirTN *size_Mat]; - D.f[dirBS ] = &DD[dirBS *size_Mat]; - D.f[dirBN ] = &DD[dirBN *size_Mat]; - D.f[dirTS ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirTNE *size_Mat]; - D.f[dirTSW ] = &DD[dirTSW *size_Mat]; - D.f[dirTSE ] = &DD[dirTSE *size_Mat]; - D.f[dirTNW ] = &DD[dirTNW *size_Mat]; - D.f[dirBNE ] = &DD[dirBNE *size_Mat]; - D.f[dirBSW ] = &DD[dirBSW *size_Mat]; - D.f[dirBSE ] = &DD[dirBSE *size_Mat]; - D.f[dirBNW ] = &DD[dirBNW *size_Mat]; - } - else - { - D.f[dirW ] = &DD[dirE *size_Mat]; - D.f[dirE ] = &DD[dirW *size_Mat]; - D.f[dirS ] = &DD[dirN *size_Mat]; - D.f[dirN ] = &DD[dirS *size_Mat]; - D.f[dirB ] = &DD[dirT *size_Mat]; - D.f[dirT ] = &DD[dirB *size_Mat]; - D.f[dirSW ] = &DD[dirNE *size_Mat]; - D.f[dirNE ] = &DD[dirSW *size_Mat]; - D.f[dirNW ] = &DD[dirSE *size_Mat]; - D.f[dirSE ] = &DD[dirNW *size_Mat]; - D.f[dirBW ] = &DD[dirTE *size_Mat]; - D.f[dirTE ] = &DD[dirBW *size_Mat]; - D.f[dirTW ] = &DD[dirBE *size_Mat]; - D.f[dirBE ] = &DD[dirTW *size_Mat]; - D.f[dirBS ] = &DD[dirTN *size_Mat]; - D.f[dirTN ] = &DD[dirBS *size_Mat]; - D.f[dirTS ] = &DD[dirBN *size_Mat]; - D.f[dirBN ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirBSW *size_Mat]; - D.f[dirTSW ] = &DD[dirBNE *size_Mat]; - D.f[dirTSE ] = &DD[dirBNW *size_Mat]; - D.f[dirTNW ] = &DD[dirBSE *size_Mat]; - D.f[dirBNE ] = &DD[dirTSW *size_Mat]; - D.f[dirBSW ] = &DD[dirTNE *size_Mat]; - D.f[dirBSE ] = &DD[dirTNW *size_Mat]; - D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } + //Distributions27 D; + //if (evenOrOdd==true) + //{ + // D.f[dirE ] = &DD[dirE *size_Mat]; + // D.f[dirW ] = &DD[dirW *size_Mat]; + // D.f[dirN ] = &DD[dirN *size_Mat]; + // D.f[dirS ] = &DD[dirS *size_Mat]; + // D.f[dirT ] = &DD[dirT *size_Mat]; + // D.f[dirB ] = &DD[dirB *size_Mat]; + // D.f[dirNE ] = &DD[dirNE *size_Mat]; + // D.f[dirSW ] = &DD[dirSW *size_Mat]; + // D.f[dirSE ] = &DD[dirSE *size_Mat]; + // D.f[dirNW ] = &DD[dirNW *size_Mat]; + // D.f[dirTE ] = &DD[dirTE *size_Mat]; + // D.f[dirBW ] = &DD[dirBW *size_Mat]; + // D.f[dirBE ] = &DD[dirBE *size_Mat]; + // D.f[dirTW ] = &DD[dirTW *size_Mat]; + // D.f[dirTN ] = &DD[dirTN *size_Mat]; + // D.f[dirBS ] = &DD[dirBS *size_Mat]; + // D.f[dirBN ] = &DD[dirBN *size_Mat]; + // D.f[dirTS ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirTNE *size_Mat]; + // D.f[dirTSW ] = &DD[dirTSW *size_Mat]; + // D.f[dirTSE ] = &DD[dirTSE *size_Mat]; + // D.f[dirTNW ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNE ] = &DD[dirBNE *size_Mat]; + // D.f[dirBSW ] = &DD[dirBSW *size_Mat]; + // D.f[dirBSE ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNW ] = &DD[dirBNW *size_Mat]; + //} + //else + //{ + // D.f[dirW ] = &DD[dirE *size_Mat]; + // D.f[dirE ] = &DD[dirW *size_Mat]; + // D.f[dirS ] = &DD[dirN *size_Mat]; + // D.f[dirN ] = &DD[dirS *size_Mat]; + // D.f[dirB ] = &DD[dirT *size_Mat]; + // D.f[dirT ] = &DD[dirB *size_Mat]; + // D.f[dirSW ] = &DD[dirNE *size_Mat]; + // D.f[dirNE ] = &DD[dirSW *size_Mat]; + // D.f[dirNW ] = &DD[dirSE *size_Mat]; + // D.f[dirSE ] = &DD[dirNW *size_Mat]; + // D.f[dirBW ] = &DD[dirTE *size_Mat]; + // D.f[dirTE ] = &DD[dirBW *size_Mat]; + // D.f[dirTW ] = &DD[dirBE *size_Mat]; + // D.f[dirBE ] = &DD[dirTW *size_Mat]; + // D.f[dirBS ] = &DD[dirTN *size_Mat]; + // D.f[dirTN ] = &DD[dirBS *size_Mat]; + // D.f[dirTS ] = &DD[dirBN *size_Mat]; + // D.f[dirBN ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirBSW *size_Mat]; + // D.f[dirTSW ] = &DD[dirBNE *size_Mat]; + // D.f[dirTSE ] = &DD[dirBNW *size_Mat]; + // D.f[dirTNW ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNE ] = &DD[dirTSW *size_Mat]; + // D.f[dirBSW ] = &DD[dirTNE *size_Mat]; + // D.f[dirBSE ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNW ] = &DD[dirTSE *size_Mat]; + //} Distributions7 D7; if (evenOrOdd==true) @@ -4059,73 +4060,73 @@ extern "C" __global__ void QNoSlipADincomp7( int inx, ////////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; + //unsigned int kzero= KQK; unsigned int ke = KQK; unsigned int kw = neighborX[KQK]; unsigned int kn = KQK; unsigned int ks = neighborY[KQK]; - unsigned int kt = KQK; - unsigned int kb = neighborZ[KQK]; - unsigned int ksw = neighborY[kw]; - unsigned int kne = KQK; - unsigned int kse = ks; - unsigned int knw = kw; - unsigned int kbw = neighborZ[kw]; - unsigned int kte = KQK; - unsigned int kbe = kb; - unsigned int ktw = kw; - unsigned int kbs = neighborZ[ks]; - unsigned int ktn = KQK; - unsigned int kbn = kb; - unsigned int kts = ks; - unsigned int ktse = ks; - unsigned int kbnw = kbw; - unsigned int ktnw = kw; - unsigned int kbse = kbs; - unsigned int ktsw = ksw; - unsigned int kbne = kb; - unsigned int ktne = KQK; - unsigned int kbsw = neighborZ[ksw]; - //////////////////////////////////////////////////////////////////////////////// - real f_W = (D.f[dirE ])[ke ]; - real f_E = (D.f[dirW ])[kw ]; - real f_S = (D.f[dirN ])[kn ]; - real f_N = (D.f[dirS ])[ks ]; - real f_B = (D.f[dirT ])[kt ]; - real f_T = (D.f[dirB ])[kb ]; - real f_SW = (D.f[dirNE ])[kne ]; - real f_NE = (D.f[dirSW ])[ksw ]; - real f_NW = (D.f[dirSE ])[kse ]; - real f_SE = (D.f[dirNW ])[knw ]; - real f_BW = (D.f[dirTE ])[kte ]; - real f_TE = (D.f[dirBW ])[kbw ]; - real f_TW = (D.f[dirBE ])[kbe ]; - real f_BE = (D.f[dirTW ])[ktw ]; - real f_BS = (D.f[dirTN ])[ktn ]; - real f_TN = (D.f[dirBS ])[kbs ]; - real f_TS = (D.f[dirBN ])[kbn ]; - real f_BN = (D.f[dirTS ])[kts ]; - real f_BSW = (D.f[dirTNE ])[ktne ]; - real f_BNE = (D.f[dirTSW ])[ktsw ]; - real f_BNW = (D.f[dirTSE ])[ktse ]; - real f_BSE = (D.f[dirTNW ])[ktnw ]; - real f_TSW = (D.f[dirBNE ])[kbne ]; - real f_TNE = (D.f[dirBSW ])[kbsw ]; - real f_TNW = (D.f[dirBSE ])[kbse ]; - real f_TSE = (D.f[dirBNW ])[kbnw ]; + unsigned int kt = KQK; + unsigned int kb = neighborZ[KQK]; + //unsigned int ksw = neighborY[kw]; + //unsigned int kne = KQK; + //unsigned int kse = ks; + //unsigned int knw = kw; + //unsigned int kbw = neighborZ[kw]; + //unsigned int kte = KQK; + //unsigned int kbe = kb; + //unsigned int ktw = kw; + //unsigned int kbs = neighborZ[ks]; + //unsigned int ktn = KQK; + //unsigned int kbn = kb; + //unsigned int kts = ks; + //unsigned int ktse = ks; + //unsigned int kbnw = kbw; + //unsigned int ktnw = kw; + //unsigned int kbse = kbs; + //unsigned int ktsw = ksw; + //unsigned int kbne = kb; + //unsigned int ktne = KQK; + //unsigned int kbsw = neighborZ[ksw]; //////////////////////////////////////////////////////////////////////////////// - real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); - real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); - real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); + //real f_W = (D.f[dirE ])[ke ]; + //real f_E = (D.f[dirW ])[kw ]; + //real f_S = (D.f[dirN ])[kn ]; + //real f_N = (D.f[dirS ])[ks ]; + //real f_B = (D.f[dirT ])[kt ]; + //real f_T = (D.f[dirB ])[kb ]; + //real f_SW = (D.f[dirNE ])[kne ]; + //real f_NE = (D.f[dirSW ])[ksw ]; + //real f_NW = (D.f[dirSE ])[kse ]; + //real f_SE = (D.f[dirNW ])[knw ]; + //real f_BW = (D.f[dirTE ])[kte ]; + //real f_TE = (D.f[dirBW ])[kbw ]; + //real f_TW = (D.f[dirBE ])[kbe ]; + //real f_BE = (D.f[dirTW ])[ktw ]; + //real f_BS = (D.f[dirTN ])[ktn ]; + //real f_TN = (D.f[dirBS ])[kbs ]; + //real f_TS = (D.f[dirBN ])[kbn ]; + //real f_BN = (D.f[dirTS ])[kts ]; + //real f_BSW = (D.f[dirTNE ])[ktne ]; + //real f_BNE = (D.f[dirTSW ])[ktsw ]; + //real f_BNW = (D.f[dirTSE ])[ktse ]; + //real f_BSE = (D.f[dirTNW ])[ktnw ]; + //real f_TSW = (D.f[dirBNE ])[kbne ]; + //real f_TNE = (D.f[dirBSW ])[kbsw ]; + //real f_TNW = (D.f[dirBSE ])[kbse ]; + //real f_TSE = (D.f[dirBNW ])[kbnw ]; + //////////////////////////////////////////////////////////////////////////////// + //real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); + //real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); + //real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); ////d�rrrrrty !!!!!!!!!!!!! // real vx1 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); // real vx2 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); // real vx3 = ten * ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //real cu_sq =c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); - real ux_sq = vx1 * vx1; - real uy_sq = vx2 * vx2; - real uz_sq = vx3 * vx3; + //real ux_sq = vx1 * vx1; + //real uy_sq = vx2 * vx2; + //real uz_sq = vx3 * vx3; //////////////////////////////////////////////////////////////////////////////// //BGK //real omegaD = three - sqrt(three); @@ -4133,13 +4134,13 @@ extern "C" __global__ void QNoSlipADincomp7( int inx, //real nue_d = Lam/three; //real ae = diffusivity/nue_d - one; //zero; - real f7_ZERO,f7_E,f7_W,f7_N,f7_S,f7_T,f7_B; + real f7_E,f7_W,f7_N,f7_S,f7_T,f7_B; //real /*feq7_ZERO,*/feq7_E,feq7_W,feq7_N,feq7_S,feq7_T,feq7_B; //real /*feqW7_ZERO,*/feqW7_E,feqW7_W,feqW7_N,feqW7_S,feqW7_T,feqW7_B; real TempD = temp[k]; - f7_ZERO = (D7.f[0])[kzero]; + //f7_ZERO = (D7.f[0])[kzero]; f7_W = (D7.f[1])[ke ]; f7_E = (D7.f[2])[kw ]; f7_S = (D7.f[3])[kn ]; @@ -4147,7 +4148,7 @@ extern "C" __global__ void QNoSlipADincomp7( int inx, f7_B = (D7.f[5])[kt ]; f7_T = (D7.f[6])[kb ]; - real ConcD = f7_ZERO + f7_E + f7_W + f7_N + f7_S + f7_T + f7_B; + //real ConcD = f7_ZERO + f7_E + f7_W + f7_N + f7_S + f7_T + f7_B; ////feq7_ZERO = ConcD*(c1o3*(ae*(-three))-(ux_sq+uy_sq+uz_sq)); //feq7_E = ConcD*(c1o6*(ae+one)+c1o2*(ux_sq)+vx1*c1o2); @@ -4859,67 +4860,67 @@ extern "C" __global__ void QADVeloIncomp7( int inx, unsigned int size_Mat, bool evenOrOdd) { - Distributions27 D; - if (evenOrOdd==true) - { - D.f[dirE ] = &DD[dirE *size_Mat]; - D.f[dirW ] = &DD[dirW *size_Mat]; - D.f[dirN ] = &DD[dirN *size_Mat]; - D.f[dirS ] = &DD[dirS *size_Mat]; - D.f[dirT ] = &DD[dirT *size_Mat]; - D.f[dirB ] = &DD[dirB *size_Mat]; - D.f[dirNE ] = &DD[dirNE *size_Mat]; - D.f[dirSW ] = &DD[dirSW *size_Mat]; - D.f[dirSE ] = &DD[dirSE *size_Mat]; - D.f[dirNW ] = &DD[dirNW *size_Mat]; - D.f[dirTE ] = &DD[dirTE *size_Mat]; - D.f[dirBW ] = &DD[dirBW *size_Mat]; - D.f[dirBE ] = &DD[dirBE *size_Mat]; - D.f[dirTW ] = &DD[dirTW *size_Mat]; - D.f[dirTN ] = &DD[dirTN *size_Mat]; - D.f[dirBS ] = &DD[dirBS *size_Mat]; - D.f[dirBN ] = &DD[dirBN *size_Mat]; - D.f[dirTS ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirTNE *size_Mat]; - D.f[dirTSW ] = &DD[dirTSW *size_Mat]; - D.f[dirTSE ] = &DD[dirTSE *size_Mat]; - D.f[dirTNW ] = &DD[dirTNW *size_Mat]; - D.f[dirBNE ] = &DD[dirBNE *size_Mat]; - D.f[dirBSW ] = &DD[dirBSW *size_Mat]; - D.f[dirBSE ] = &DD[dirBSE *size_Mat]; - D.f[dirBNW ] = &DD[dirBNW *size_Mat]; - } - else - { - D.f[dirW ] = &DD[dirE *size_Mat]; - D.f[dirE ] = &DD[dirW *size_Mat]; - D.f[dirS ] = &DD[dirN *size_Mat]; - D.f[dirN ] = &DD[dirS *size_Mat]; - D.f[dirB ] = &DD[dirT *size_Mat]; - D.f[dirT ] = &DD[dirB *size_Mat]; - D.f[dirSW ] = &DD[dirNE *size_Mat]; - D.f[dirNE ] = &DD[dirSW *size_Mat]; - D.f[dirNW ] = &DD[dirSE *size_Mat]; - D.f[dirSE ] = &DD[dirNW *size_Mat]; - D.f[dirBW ] = &DD[dirTE *size_Mat]; - D.f[dirTE ] = &DD[dirBW *size_Mat]; - D.f[dirTW ] = &DD[dirBE *size_Mat]; - D.f[dirBE ] = &DD[dirTW *size_Mat]; - D.f[dirBS ] = &DD[dirTN *size_Mat]; - D.f[dirTN ] = &DD[dirBS *size_Mat]; - D.f[dirTS ] = &DD[dirBN *size_Mat]; - D.f[dirBN ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirBSW *size_Mat]; - D.f[dirTSW ] = &DD[dirBNE *size_Mat]; - D.f[dirTSE ] = &DD[dirBNW *size_Mat]; - D.f[dirTNW ] = &DD[dirBSE *size_Mat]; - D.f[dirBNE ] = &DD[dirTSW *size_Mat]; - D.f[dirBSW ] = &DD[dirTNE *size_Mat]; - D.f[dirBSE ] = &DD[dirTNW *size_Mat]; - D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } + //Distributions27 D; + //if (evenOrOdd==true) + //{ + // D.f[dirE ] = &DD[dirE *size_Mat]; + // D.f[dirW ] = &DD[dirW *size_Mat]; + // D.f[dirN ] = &DD[dirN *size_Mat]; + // D.f[dirS ] = &DD[dirS *size_Mat]; + // D.f[dirT ] = &DD[dirT *size_Mat]; + // D.f[dirB ] = &DD[dirB *size_Mat]; + // D.f[dirNE ] = &DD[dirNE *size_Mat]; + // D.f[dirSW ] = &DD[dirSW *size_Mat]; + // D.f[dirSE ] = &DD[dirSE *size_Mat]; + // D.f[dirNW ] = &DD[dirNW *size_Mat]; + // D.f[dirTE ] = &DD[dirTE *size_Mat]; + // D.f[dirBW ] = &DD[dirBW *size_Mat]; + // D.f[dirBE ] = &DD[dirBE *size_Mat]; + // D.f[dirTW ] = &DD[dirTW *size_Mat]; + // D.f[dirTN ] = &DD[dirTN *size_Mat]; + // D.f[dirBS ] = &DD[dirBS *size_Mat]; + // D.f[dirBN ] = &DD[dirBN *size_Mat]; + // D.f[dirTS ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirTNE *size_Mat]; + // D.f[dirTSW ] = &DD[dirTSW *size_Mat]; + // D.f[dirTSE ] = &DD[dirTSE *size_Mat]; + // D.f[dirTNW ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNE ] = &DD[dirBNE *size_Mat]; + // D.f[dirBSW ] = &DD[dirBSW *size_Mat]; + // D.f[dirBSE ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNW ] = &DD[dirBNW *size_Mat]; + //} + //else + //{ + // D.f[dirW ] = &DD[dirE *size_Mat]; + // D.f[dirE ] = &DD[dirW *size_Mat]; + // D.f[dirS ] = &DD[dirN *size_Mat]; + // D.f[dirN ] = &DD[dirS *size_Mat]; + // D.f[dirB ] = &DD[dirT *size_Mat]; + // D.f[dirT ] = &DD[dirB *size_Mat]; + // D.f[dirSW ] = &DD[dirNE *size_Mat]; + // D.f[dirNE ] = &DD[dirSW *size_Mat]; + // D.f[dirNW ] = &DD[dirSE *size_Mat]; + // D.f[dirSE ] = &DD[dirNW *size_Mat]; + // D.f[dirBW ] = &DD[dirTE *size_Mat]; + // D.f[dirTE ] = &DD[dirBW *size_Mat]; + // D.f[dirTW ] = &DD[dirBE *size_Mat]; + // D.f[dirBE ] = &DD[dirTW *size_Mat]; + // D.f[dirBS ] = &DD[dirTN *size_Mat]; + // D.f[dirTN ] = &DD[dirBS *size_Mat]; + // D.f[dirTS ] = &DD[dirBN *size_Mat]; + // D.f[dirBN ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirBSW *size_Mat]; + // D.f[dirTSW ] = &DD[dirBNE *size_Mat]; + // D.f[dirTSE ] = &DD[dirBNW *size_Mat]; + // D.f[dirTNW ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNE ] = &DD[dirTSW *size_Mat]; + // D.f[dirBSW ] = &DD[dirTNE *size_Mat]; + // D.f[dirBSE ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNW ] = &DD[dirTSE *size_Mat]; + //} Distributions7 D7; if (evenOrOdd==true) @@ -4969,87 +4970,87 @@ extern "C" __global__ void QADVeloIncomp7( int inx, ////////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; + //unsigned int kzero= KQK; unsigned int ke = KQK; unsigned int kw = neighborX[KQK]; unsigned int kn = KQK; unsigned int ks = neighborY[KQK]; unsigned int kt = KQK; unsigned int kb = neighborZ[KQK]; - unsigned int ksw = neighborY[kw]; - unsigned int kne = KQK; - unsigned int kse = ks; - unsigned int knw = kw; - unsigned int kbw = neighborZ[kw]; - unsigned int kte = KQK; - unsigned int kbe = kb; - unsigned int ktw = kw; - unsigned int kbs = neighborZ[ks]; - unsigned int ktn = KQK; - unsigned int kbn = kb; - unsigned int kts = ks; - unsigned int ktse = ks; - unsigned int kbnw = kbw; - unsigned int ktnw = kw; - unsigned int kbse = kbs; - unsigned int ktsw = ksw; - unsigned int kbne = kb; - unsigned int ktne = KQK; - unsigned int kbsw = neighborZ[ksw]; + //unsigned int ksw = neighborY[kw]; + //unsigned int kne = KQK; + //unsigned int kse = ks; + //unsigned int knw = kw; + //unsigned int kbw = neighborZ[kw]; + //unsigned int kte = KQK; + //unsigned int kbe = kb; + //unsigned int ktw = kw; + //unsigned int kbs = neighborZ[ks]; + //unsigned int ktn = KQK; + //unsigned int kbn = kb; + //unsigned int kts = ks; + //unsigned int ktse = ks; + //unsigned int kbnw = kbw; + //unsigned int ktnw = kw; + //unsigned int kbse = kbs; + //unsigned int ktsw = ksw; + //unsigned int kbne = kb; + //unsigned int ktne = KQK; + //unsigned int kbsw = neighborZ[ksw]; //////////////////////////////////////////////////////////////////////////////// - real f_W = (D.f[dirE ])[ke ]; - real f_E = (D.f[dirW ])[kw ]; - real f_S = (D.f[dirN ])[kn ]; - real f_N = (D.f[dirS ])[ks ]; - real f_B = (D.f[dirT ])[kt ]; - real f_T = (D.f[dirB ])[kb ]; - real f_SW = (D.f[dirNE ])[kne ]; - real f_NE = (D.f[dirSW ])[ksw ]; - real f_NW = (D.f[dirSE ])[kse ]; - real f_SE = (D.f[dirNW ])[knw ]; - real f_BW = (D.f[dirTE ])[kte ]; - real f_TE = (D.f[dirBW ])[kbw ]; - real f_TW = (D.f[dirBE ])[kbe ]; - real f_BE = (D.f[dirTW ])[ktw ]; - real f_BS = (D.f[dirTN ])[ktn ]; - real f_TN = (D.f[dirBS ])[kbs ]; - real f_TS = (D.f[dirBN ])[kbn ]; - real f_BN = (D.f[dirTS ])[kts ]; - real f_BSW = (D.f[dirTNE ])[ktne ]; - real f_BNE = (D.f[dirTSW ])[ktsw ]; - real f_BNW = (D.f[dirTSE ])[ktse ]; - real f_BSE = (D.f[dirTNW ])[ktnw ]; - real f_TSW = (D.f[dirBNE ])[kbne ]; - real f_TNE = (D.f[dirBSW ])[kbsw ]; - real f_TNW = (D.f[dirBSE ])[kbse ]; - real f_TSE = (D.f[dirBNW ])[kbnw ]; + //real f_W = (D.f[dirE ])[ke ]; + //real f_E = (D.f[dirW ])[kw ]; + //real f_S = (D.f[dirN ])[kn ]; + //real f_N = (D.f[dirS ])[ks ]; + //real f_B = (D.f[dirT ])[kt ]; + //real f_T = (D.f[dirB ])[kb ]; + //real f_SW = (D.f[dirNE ])[kne ]; + //real f_NE = (D.f[dirSW ])[ksw ]; + //real f_NW = (D.f[dirSE ])[kse ]; + //real f_SE = (D.f[dirNW ])[knw ]; + //real f_BW = (D.f[dirTE ])[kte ]; + //real f_TE = (D.f[dirBW ])[kbw ]; + //real f_TW = (D.f[dirBE ])[kbe ]; + //real f_BE = (D.f[dirTW ])[ktw ]; + //real f_BS = (D.f[dirTN ])[ktn ]; + //real f_TN = (D.f[dirBS ])[kbs ]; + //real f_TS = (D.f[dirBN ])[kbn ]; + //real f_BN = (D.f[dirTS ])[kts ]; + //real f_BSW = (D.f[dirTNE ])[ktne ]; + //real f_BNE = (D.f[dirTSW ])[ktsw ]; + //real f_BNW = (D.f[dirTSE ])[ktse ]; + //real f_BSE = (D.f[dirTNW ])[ktnw ]; + //real f_TSW = (D.f[dirBNE ])[kbne ]; + //real f_TNE = (D.f[dirBSW ])[kbsw ]; + //real f_TNW = (D.f[dirBSE ])[kbse ]; + //real f_TSE = (D.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// - real vx1_Inflow = c0o1; - real vx2_Inflow = velo[k]; - real vx3_Inflow = c0o1; - real ux_sq_Inflow = vx1_Inflow * vx1_Inflow; - real uy_sq_Inflow = vx2_Inflow * vx2_Inflow; - real uz_sq_Inflow = vx3_Inflow * vx3_Inflow; + //real vx1_Inflow = c0o1; + //real vx2_Inflow = velo[k]; + //real vx3_Inflow = c0o1; + //real ux_sq_Inflow = vx1_Inflow * vx1_Inflow; + //real uy_sq_Inflow = vx2_Inflow * vx2_Inflow; + //real uz_sq_Inflow = vx3_Inflow * vx3_Inflow; //////////////////////////////////////////////////////////////////////////////// - real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); - real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); - real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); + //real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); + //real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); + //real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); ////d�rrrrrty !!!!!!!!!!!!! // real vx1 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); // real vx2 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); // real vx3 = ten * ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //////////////////////////////////////////////////////////////////////////////// //real cu_sq =1.5f*(vx1*vx1+vx2*vx2+vx3*vx3); - real ux_sq = vx1 * vx1; - real uy_sq = vx2 * vx2; - real uz_sq = vx3 * vx3; + //real ux_sq = vx1 * vx1; + //real uy_sq = vx2 * vx2; + //real uz_sq = vx3 * vx3; - real f7_ZERO,f7_E,f7_W,f7_N,f7_S,f7_T,f7_B; + real f7_E,f7_W,f7_N,f7_S,f7_T,f7_B; //real /*feq7_ZERO,*/feq7_E,feq7_W,feq7_N,feq7_S,feq7_T,feq7_B; //real /*feqW7_ZERO,*/feqW7_E,feqW7_W,feqW7_N,feqW7_S,feqW7_T,feqW7_B; real TempD = temp[k]; - f7_ZERO = (D7.f[0])[kzero]; + //f7_ZERO = (D7.f[0])[kzero]; f7_W = (D7.f[1])[ke ]; f7_E = (D7.f[2])[kw ]; f7_S = (D7.f[3])[kn ]; @@ -5057,7 +5058,7 @@ extern "C" __global__ void QADVeloIncomp7( int inx, f7_B = (D7.f[5])[kt ]; f7_T = (D7.f[6])[kb ]; - real ConcD = f7_ZERO + f7_E + f7_W + f7_N + f7_S + f7_T + f7_B; + //real ConcD = f7_ZERO + f7_E + f7_W + f7_N + f7_S + f7_T + f7_B; //////////////////////////////////////////////////////////////////////////////// //BGK @@ -5510,7 +5511,7 @@ extern "C" __global__ void QADVeloIncomp27( int inx, //////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; + //unsigned int kzero= KQK; unsigned int ke = KQK; unsigned int kw = neighborX[KQK]; unsigned int kn = KQK; @@ -5570,67 +5571,67 @@ extern "C" __global__ void QADVeloIncomp27( int inx, real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //////////////////////////////////////////////////////////////////////////////// - real f27_W = (D27.f[dirE ])[ke ]; - real f27_E = (D27.f[dirW ])[kw ]; - real f27_S = (D27.f[dirN ])[kn ]; - real f27_N = (D27.f[dirS ])[ks ]; - real f27_B = (D27.f[dirT ])[kt ]; - real f27_T = (D27.f[dirB ])[kb ]; - real f27_SW = (D27.f[dirNE ])[kne ]; - real f27_NE = (D27.f[dirSW ])[ksw ]; - real f27_NW = (D27.f[dirSE ])[kse ]; - real f27_SE = (D27.f[dirNW ])[knw ]; - real f27_BW = (D27.f[dirTE ])[kte ]; - real f27_TE = (D27.f[dirBW ])[kbw ]; - real f27_TW = (D27.f[dirBE ])[kbe ]; - real f27_BE = (D27.f[dirTW ])[ktw ]; - real f27_BS = (D27.f[dirTN ])[ktn ]; - real f27_TN = (D27.f[dirBS ])[kbs ]; - real f27_TS = (D27.f[dirBN ])[kbn ]; - real f27_BN = (D27.f[dirTS ])[kts ]; - real f27_ZERO = (D27.f[dirZERO])[kzero]; - real f27_BSW = (D27.f[dirTNE ])[ktne ]; - real f27_BNE = (D27.f[dirTSW ])[ktsw ]; - real f27_BNW = (D27.f[dirTSE ])[ktse ]; - real f27_BSE = (D27.f[dirTNW ])[ktnw ]; - real f27_TSW = (D27.f[dirBNE ])[kbne ]; - real f27_TNE = (D27.f[dirBSW ])[kbsw ]; - real f27_TNW = (D27.f[dirBSE ])[kbse ]; - real f27_TSE = (D27.f[dirBNW ])[kbnw ]; + //real f27_W = (D27.f[dirE ])[ke ]; + //real f27_E = (D27.f[dirW ])[kw ]; + //real f27_S = (D27.f[dirN ])[kn ]; + //real f27_N = (D27.f[dirS ])[ks ]; + //real f27_B = (D27.f[dirT ])[kt ]; + //real f27_T = (D27.f[dirB ])[kb ]; + //real f27_SW = (D27.f[dirNE ])[kne ]; + //real f27_NE = (D27.f[dirSW ])[ksw ]; + //real f27_NW = (D27.f[dirSE ])[kse ]; + //real f27_SE = (D27.f[dirNW ])[knw ]; + //real f27_BW = (D27.f[dirTE ])[kte ]; + //real f27_TE = (D27.f[dirBW ])[kbw ]; + //real f27_TW = (D27.f[dirBE ])[kbe ]; + //real f27_BE = (D27.f[dirTW ])[ktw ]; + //real f27_BS = (D27.f[dirTN ])[ktn ]; + //real f27_TN = (D27.f[dirBS ])[kbs ]; + //real f27_TS = (D27.f[dirBN ])[kbn ]; + //real f27_BN = (D27.f[dirTS ])[kts ]; + //real f27_ZERO = (D27.f[dirZERO])[kzero]; + //real f27_BSW = (D27.f[dirTNE ])[ktne ]; + //real f27_BNE = (D27.f[dirTSW ])[ktsw ]; + //real f27_BNW = (D27.f[dirTSE ])[ktse ]; + //real f27_BSE = (D27.f[dirTNW ])[ktnw ]; + //real f27_TSW = (D27.f[dirBNE ])[kbne ]; + //real f27_TNE = (D27.f[dirBSW ])[kbsw ]; + //real f27_TNW = (D27.f[dirBSE ])[kbse ]; + //real f27_TSE = (D27.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); //////////////////////////////////////////////////////////////////////////////// - real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + - f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + - f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; + //real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + + // f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + + // f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; //real feq27_ZERO = c8over27* ConcD*(one-cu_sq); - real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); - real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); - real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); - real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); - real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); - real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); - real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); - real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); - real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); - real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); - real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); - real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); - real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); - real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); - real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); - real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); - real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); - real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); - real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); - real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); - real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); - real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); - real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); - real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); - real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); + //real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); + //real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); + //real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); + //real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); + //real feq27_T = c2o27* ConcD*(c1o1+c3o1*( vx3)+c9o2*( vx3)*( vx3)-cu_sq); + //real feq27_B = c2o27* ConcD*(c1o1+c3o1*( -vx3)+c9o2*( -vx3)*( -vx3)-cu_sq); + //real feq27_NE = c1o54* ConcD*(c1o1+c3o1*( vx1+vx2 )+c9o2*( vx1+vx2 )*( vx1+vx2 )-cu_sq); + //real feq27_SW = c1o54* ConcD*(c1o1+c3o1*(-vx1-vx2 )+c9o2*(-vx1-vx2 )*(-vx1-vx2 )-cu_sq); + //real feq27_SE = c1o54* ConcD*(c1o1+c3o1*( vx1-vx2 )+c9o2*( vx1-vx2 )*( vx1-vx2 )-cu_sq); + //real feq27_NW = c1o54* ConcD*(c1o1+c3o1*(-vx1+vx2 )+c9o2*(-vx1+vx2 )*(-vx1+vx2 )-cu_sq); + //real feq27_TE = c1o54* ConcD*(c1o1+c3o1*( vx1 +vx3)+c9o2*( vx1 +vx3)*( vx1 +vx3)-cu_sq); + //real feq27_BW = c1o54* ConcD*(c1o1+c3o1*(-vx1 -vx3)+c9o2*(-vx1 -vx3)*(-vx1 -vx3)-cu_sq); + //real feq27_BE = c1o54* ConcD*(c1o1+c3o1*( vx1 -vx3)+c9o2*( vx1 -vx3)*( vx1 -vx3)-cu_sq); + //real feq27_TW = c1o54* ConcD*(c1o1+c3o1*(-vx1 +vx3)+c9o2*(-vx1 +vx3)*(-vx1 +vx3)-cu_sq); + //real feq27_TN = c1o54* ConcD*(c1o1+c3o1*( vx2+vx3)+c9o2*( vx2+vx3)*( vx2+vx3)-cu_sq); + //real feq27_BS = c1o54* ConcD*(c1o1+c3o1*( -vx2-vx3)+c9o2*( -vx2-vx3)*( -vx2-vx3)-cu_sq); + //real feq27_BN = c1o54* ConcD*(c1o1+c3o1*( vx2-vx3)+c9o2*( vx2-vx3)*( vx2-vx3)-cu_sq); + //real feq27_TS = c1o54* ConcD*(c1o1+c3o1*( -vx2+vx3)+c9o2*( -vx2+vx3)*( -vx2+vx3)-cu_sq); + //real feq27_TNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2+vx3)+c9o2*( vx1+vx2+vx3)*( vx1+vx2+vx3)-cu_sq); + //real feq27_BSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2-vx3)+c9o2*(-vx1-vx2-vx3)*(-vx1-vx2-vx3)-cu_sq); + //real feq27_BNE = c1o216*ConcD*(c1o1+c3o1*( vx1+vx2-vx3)+c9o2*( vx1+vx2-vx3)*( vx1+vx2-vx3)-cu_sq); + //real feq27_TSW = c1o216*ConcD*(c1o1+c3o1*(-vx1-vx2+vx3)+c9o2*(-vx1-vx2+vx3)*(-vx1-vx2+vx3)-cu_sq); + //real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); + //real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); + //real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); + //real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// real TempD = temp[k]; //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -5854,7 +5855,7 @@ extern "C" __global__ void QADPressIncomp7(int inx, unsigned int size_Mat, bool evenOrOdd) { - Distributions27 D; + /* Distributions27 D; if (evenOrOdd==true) { D.f[dirE ] = &DD[dirE *size_Mat]; @@ -5914,7 +5915,7 @@ extern "C" __global__ void QADPressIncomp7(int inx, D.f[dirBSW ] = &DD[dirTNE *size_Mat]; D.f[dirBSE ] = &DD[dirTNW *size_Mat]; D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } + }*/ Distributions7 D7; if (evenOrOdd==true) @@ -5971,28 +5972,28 @@ extern "C" __global__ void QADPressIncomp7(int inx, unsigned int ks = neighborY[KQK]; unsigned int kt = KQK; unsigned int kb = neighborZ[KQK]; - unsigned int ksw = neighborY[kw]; - unsigned int kne = KQK; - unsigned int kse = ks; - unsigned int knw = kw; - unsigned int kbw = neighborZ[kw]; - unsigned int kte = KQK; - unsigned int kbe = kb; - unsigned int ktw = kw; - unsigned int kbs = neighborZ[ks]; - unsigned int ktn = KQK; - unsigned int kbn = kb; - unsigned int kts = ks; - unsigned int ktse = ks; - unsigned int kbnw = kbw; - unsigned int ktnw = kw; - unsigned int kbse = kbs; - unsigned int ktsw = ksw; - unsigned int kbne = kb; - unsigned int ktne = KQK; - unsigned int kbsw = neighborZ[ksw]; + //unsigned int ksw = neighborY[kw]; + //unsigned int kne = KQK; + //unsigned int kse = ks; + //unsigned int knw = kw; + //unsigned int kbw = neighborZ[kw]; + //unsigned int kte = KQK; + //unsigned int kbe = kb; + //unsigned int ktw = kw; + //unsigned int kbs = neighborZ[ks]; + //unsigned int ktn = KQK; + //unsigned int kbn = kb; + //unsigned int kts = ks; + //unsigned int ktse = ks; + //unsigned int kbnw = kbw; + //unsigned int ktnw = kw; + //unsigned int kbse = kbs; + //unsigned int ktsw = ksw; + //unsigned int kbne = kb; + //unsigned int ktne = KQK; + //unsigned int kbsw = neighborZ[ksw]; //////////////////////////////////////////////////////////////////////////////// - real f_E, f_W, f_N, f_S, f_T, f_B, f_NE, f_SW, f_SE, f_NW, f_TE, f_BW, f_BE, + /* real f_E, f_W, f_N, f_S, f_T, f_B, f_NE, f_SW, f_SE, f_NW, f_TE, f_BW, f_BE, f_TW, f_TN, f_BS, f_BN, f_TS, f_TNE, f_TSW, f_TSE, f_TNW, f_BNE, f_BSW, f_BSE, f_BNW; f_W = (D.f[dirE ])[ke ]; @@ -6020,20 +6021,20 @@ extern "C" __global__ void QADPressIncomp7(int inx, f_TSW = (D.f[dirBNE ])[kbne ]; f_TNE = (D.f[dirBSW ])[kbsw ]; f_TNW = (D.f[dirBSE ])[kbse ]; - f_TSE = (D.f[dirBNW ])[kbnw ]; + f_TSE = (D.f[dirBNW ])[kbnw ];*/ //////////////////////////////////////////////////////////////////////////////// - real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); - real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); - real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); + //real vx1 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); + //real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); + //real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); ////d�rrrrrty !!!!!!!!!!!!! // real vx1 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_TSE-f_BNW)+(f_BSE-f_TNW) +(f_NE-f_SW)+(f_SE-f_NW)+(f_TE-f_BW)+(f_BE-f_TW)+(f_E-f_W)); // real vx2 = ten * ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); // real vx3 = ten * ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //real cu_sq =1.5*(vx1*vx1+vx2*vx2+vx3*vx3); - real ux_sq = vx1 * vx1; - real uy_sq = vx2 * vx2; - real uz_sq = vx3 * vx3; + //real ux_sq = vx1 * vx1; + //real uy_sq = vx2 * vx2; + //real uz_sq = vx3 * vx3; ////////////////////////////////////////////////////////////////////////// //BGK //real omegaD = three - sqrt(three); @@ -6078,7 +6079,7 @@ extern "C" __global__ void QADPressIncomp7(int inx, //TRT Yoshida Kernel - based on Ying real cs2 = c1o4; real Lam = diffusivity/(c1o1)/cs2; - real omegaD = - c1o1 / (Lam + c1o2); + //real omegaD = - c1o1 / (Lam + c1o2); real nue_d = Lam/c3o1; ////////////////////////////////////////////////////////////////////////// @@ -6466,7 +6467,7 @@ extern "C" __global__ void QADPressIncomp27( int inx, //////////////////////////////////////////////////////////////////////////////// //index unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; + //unsigned int kzero= KQK; unsigned int ke = KQK; unsigned int kw = neighborX[KQK]; unsigned int kn = KQK; @@ -6526,42 +6527,42 @@ extern "C" __global__ void QADPressIncomp27( int inx, real vx2 = ((f_TNE-f_BSW)+(f_BNE-f_TSW)+(f_BNW-f_TSE)+(f_TNW-f_BSE) +(f_NE-f_SW)+(f_NW-f_SE)+(f_TN-f_BS)+(f_BN-f_TS)+(f_N-f_S)); real vx3 = ((f_TNE-f_BSW)+(f_TSW-f_BNE)+(f_TSE-f_BNW)+(f_TNW-f_BSE) +(f_TE-f_BW)+(f_TW-f_BE)+(f_TN-f_BS)+(f_TS-f_BN)+(f_T-f_B)); //////////////////////////////////////////////////////////////////////////////// - real f27_W = (D27.f[dirE ])[ke ]; - real f27_E = (D27.f[dirW ])[kw ]; - real f27_S = (D27.f[dirN ])[kn ]; - real f27_N = (D27.f[dirS ])[ks ]; - real f27_B = (D27.f[dirT ])[kt ]; - real f27_T = (D27.f[dirB ])[kb ]; - real f27_SW = (D27.f[dirNE ])[kne ]; - real f27_NE = (D27.f[dirSW ])[ksw ]; - real f27_NW = (D27.f[dirSE ])[kse ]; - real f27_SE = (D27.f[dirNW ])[knw ]; - real f27_BW = (D27.f[dirTE ])[kte ]; - real f27_TE = (D27.f[dirBW ])[kbw ]; - real f27_TW = (D27.f[dirBE ])[kbe ]; - real f27_BE = (D27.f[dirTW ])[ktw ]; - real f27_BS = (D27.f[dirTN ])[ktn ]; - real f27_TN = (D27.f[dirBS ])[kbs ]; - real f27_TS = (D27.f[dirBN ])[kbn ]; - real f27_BN = (D27.f[dirTS ])[kts ]; - real f27_ZERO = (D27.f[dirZERO])[kzero]; - real f27_BSW = (D27.f[dirTNE ])[ktne ]; - real f27_BNE = (D27.f[dirTSW ])[ktsw ]; - real f27_BNW = (D27.f[dirTSE ])[ktse ]; - real f27_BSE = (D27.f[dirTNW ])[ktnw ]; - real f27_TSW = (D27.f[dirBNE ])[kbne ]; - real f27_TNE = (D27.f[dirBSW ])[kbsw ]; - real f27_TNW = (D27.f[dirBSE ])[kbse ]; - real f27_TSE = (D27.f[dirBNW ])[kbnw ]; + //real f27_W = (D27.f[dirE ])[ke ]; + //real f27_E = (D27.f[dirW ])[kw ]; + //real f27_S = (D27.f[dirN ])[kn ]; + //real f27_N = (D27.f[dirS ])[ks ]; + //real f27_B = (D27.f[dirT ])[kt ]; + //real f27_T = (D27.f[dirB ])[kb ]; + //real f27_SW = (D27.f[dirNE ])[kne ]; + //real f27_NE = (D27.f[dirSW ])[ksw ]; + //real f27_NW = (D27.f[dirSE ])[kse ]; + //real f27_SE = (D27.f[dirNW ])[knw ]; + //real f27_BW = (D27.f[dirTE ])[kte ]; + //real f27_TE = (D27.f[dirBW ])[kbw ]; + //real f27_TW = (D27.f[dirBE ])[kbe ]; + //real f27_BE = (D27.f[dirTW ])[ktw ]; + //real f27_BS = (D27.f[dirTN ])[ktn ]; + //real f27_TN = (D27.f[dirBS ])[kbs ]; + //real f27_TS = (D27.f[dirBN ])[kbn ]; + //real f27_BN = (D27.f[dirTS ])[kts ]; + //real f27_ZERO = (D27.f[dirZERO])[kzero]; + //real f27_BSW = (D27.f[dirTNE ])[ktne ]; + //real f27_BNE = (D27.f[dirTSW ])[ktsw ]; + //real f27_BNW = (D27.f[dirTSE ])[ktse ]; + //real f27_BSE = (D27.f[dirTNW ])[ktnw ]; + //real f27_TSW = (D27.f[dirBNE ])[kbne ]; + //real f27_TNE = (D27.f[dirBSW ])[kbsw ]; + //real f27_TNW = (D27.f[dirBSE ])[kbse ]; + //real f27_TSE = (D27.f[dirBNW ])[kbnw ]; //////////////////////////////////////////////////////////////////////////////// real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3); //////////////////////////////////////////////////////////////////////////////// - real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + - f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + - f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; + //real ConcD = f27_TSE + f27_TNW + f27_TNE + f27_TSW + f27_BSE + f27_BNW + f27_BNE + f27_BSW + + //f27_BN + f27_TS + f27_TN + f27_BS + f27_BE + f27_TW + f27_TE + f27_BW + f27_SE + f27_NW + f27_NE + f27_SW + + //f27_T + f27_B + f27_N + f27_S + f27_E + f27_W + f27_ZERO; //real feq27_ZERO = c8over27* ConcD*(one-cu_sq); - real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); + /*real feq27_E = c2o27* ConcD*(c1o1+c3o1*( vx1 )+c9o2*( vx1 )*( vx1 )-cu_sq); real feq27_W = c2o27* ConcD*(c1o1+c3o1*(-vx1 )+c9o2*(-vx1 )*(-vx1 )-cu_sq); real feq27_N = c2o27* ConcD*(c1o1+c3o1*( vx2 )+c9o2*( vx2 )*( vx2 )-cu_sq); real feq27_S = c2o27* ConcD*(c1o1+c3o1*( -vx2 )+c9o2*( -vx2 )*( -vx2 )-cu_sq); @@ -6586,7 +6587,7 @@ extern "C" __global__ void QADPressIncomp27( int inx, real feq27_TSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2+vx3)+c9o2*( vx1-vx2+vx3)*( vx1-vx2+vx3)-cu_sq); real feq27_BNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2-vx3)+c9o2*(-vx1+vx2-vx3)*(-vx1+vx2-vx3)-cu_sq); real feq27_BSE = c1o216*ConcD*(c1o1+c3o1*( vx1-vx2-vx3)+c9o2*( vx1-vx2-vx3)*( vx1-vx2-vx3)-cu_sq); - real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq); + real feq27_TNW = c1o216*ConcD*(c1o1+c3o1*(-vx1+vx2+vx3)+c9o2*(-vx1+vx2+vx3)*(-vx1+vx2+vx3)-cu_sq);*/ //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// real TempD = temp[k]; //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/VirtualFluids_GPU/GPU/Cumulant_F3_27.cu b/src/gpu/VirtualFluids_GPU/GPU/Cumulant_F3_27.cu index feb392b1d..035384cdf 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/Cumulant_F3_27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/Cumulant_F3_27.cu @@ -478,33 +478,33 @@ extern "C" __global__ void LB_PostProcessor_F3_2018_Fehlberg(real omega, //////////////////////////////////////////////////////////// //central moments to cumulants //4. - real CUMcbb = mfcbb - ((mfcaa + c1o3) * mfabb + c2o1 * mfbba * mfbab) / rho; - real CUMbcb = mfbcb - ((mfaca + c1o3) * mfbab + c2o1 * mfbba * mfabb) / rho; - real CUMbbc = mfbbc - ((mfaac + c1o3) * mfbba + c2o1 * mfbab * mfabb) / rho; + //real CUMcbb = mfcbb - ((mfcaa + c1o3) * mfabb + c2o1 * mfbba * mfbab) / rho; + //real CUMbcb = mfbcb - ((mfaca + c1o3) * mfbab + c2o1 * mfbba * mfabb) / rho; + //real CUMbbc = mfbbc - ((mfaac + c1o3) * mfbba + c2o1 * mfbab * mfabb) / rho; - real CUMcca = mfcca - (((mfcaa * mfaca + c2o1 * mfbba * mfbba) + c1o3 * (mfcaa + mfaca)) / rho - c1o9*(drho / rho)); - real CUMcac = mfcac - (((mfcaa * mfaac + c2o1 * mfbab * mfbab) + c1o3 * (mfcaa + mfaac)) / rho - c1o9*(drho / rho)); - real CUMacc = mfacc - (((mfaac * mfaca + c2o1 * mfabb * mfabb) + c1o3 * (mfaac + mfaca)) / rho - c1o9*(drho / rho)); + //real CUMcca = mfcca - (((mfcaa * mfaca + c2o1 * mfbba * mfbba) + c1o3 * (mfcaa + mfaca)) / rho - c1o9*(drho / rho)); + //real CUMcac = mfcac - (((mfcaa * mfaac + c2o1 * mfbab * mfbab) + c1o3 * (mfcaa + mfaac)) / rho - c1o9*(drho / rho)); + //real CUMacc = mfacc - (((mfaac * mfaca + c2o1 * mfabb * mfabb) + c1o3 * (mfaac + mfaca)) / rho - c1o9*(drho / rho)); //5. - real CUMbcc = mfbcc - ((mfaac * mfbca + mfaca * mfbac + c4o1 * mfabb * mfbbb + c2o1 * (mfbab * mfacb + mfbba * mfabc)) + c1o3 * (mfbca + mfbac)) / rho; - real CUMcbc = mfcbc - ((mfaac * mfcba + mfcaa * mfabc + c4o1 * mfbab * mfbbb + c2o1 * (mfabb * mfcab + mfbba * mfbac)) + c1o3 * (mfcba + mfabc)) / rho; - real CUMccb = mfccb - ((mfcaa * mfacb + mfaca * mfcab + c4o1 * mfbba * mfbbb + c2o1 * (mfbab * mfbca + mfabb * mfcba)) + c1o3 * (mfacb + mfcab)) / rho; + //real CUMbcc = mfbcc - ((mfaac * mfbca + mfaca * mfbac + c4o1 * mfabb * mfbbb + c2o1 * (mfbab * mfacb + mfbba * mfabc)) + c1o3 * (mfbca + mfbac)) / rho; + //real CUMcbc = mfcbc - ((mfaac * mfcba + mfcaa * mfabc + c4o1 * mfbab * mfbbb + c2o1 * (mfabb * mfcab + mfbba * mfbac)) + c1o3 * (mfcba + mfabc)) / rho; + //real CUMccb = mfccb - ((mfcaa * mfacb + mfaca * mfcab + c4o1 * mfbba * mfbbb + c2o1 * (mfbab * mfbca + mfabb * mfcba)) + c1o3 * (mfacb + mfcab)) / rho; //6. - real CUMccc = mfccc + ((-c4o1 * mfbbb * mfbbb - - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca) - - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc) - - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) / rho - + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac) - + c2o1 * (mfcaa * mfaca * mfaac) - + c16o1 * mfbba * mfbab * mfabb) / (rho * rho) - - c1o3 * (mfacc + mfcac + mfcca) / rho - - c1o9 * (mfcaa + mfaca + mfaac) / rho - + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba) - + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) / (rho * rho) * c2o3 - + c1o27*((drho * drho - drho) / (rho*rho))); + //real CUMccc = mfccc + ((-c4o1 * mfbbb * mfbbb + // - (mfcaa * mfacc + mfaca * mfcac + mfaac * mfcca) + // - c4o1 * (mfabb * mfcbb + mfbab * mfbcb + mfbba * mfbbc) + // - c2o1 * (mfbca * mfbac + mfcba * mfabc + mfcab * mfacb)) / rho + // + (c4o1 * (mfbab * mfbab * mfaca + mfabb * mfabb * mfcaa + mfbba * mfbba * mfaac) + // + c2o1 * (mfcaa * mfaca * mfaac) + // + c16o1 * mfbba * mfbab * mfabb) / (rho * rho) + // - c1o3 * (mfacc + mfcac + mfcca) / rho + // - c1o9 * (mfcaa + mfaca + mfaac) / rho + // + (c2o1 * (mfbab * mfbab + mfabb * mfabb + mfbba * mfbba) + // + (mfaac * mfaca + mfaac * mfcaa + mfaca * mfcaa) + c1o3 *(mfaac + mfaca + mfcaa)) / (rho * rho) * c2o3 + // + c1o27*((drho * drho - drho) / (rho*rho))); //2. // linear combinations @@ -513,9 +513,9 @@ extern "C" __global__ void LB_PostProcessor_F3_2018_Fehlberg(real omega, real mxxMzz = mfcaa - mfaac; //////////////////////////////////////////////////////////////////////////// - real Dxy = -c3o1*omega*mfbba; - real Dxz = -c3o1*omega*mfbab; - real Dyz = -c3o1*omega*mfabb; + //real Dxy = -c3o1*omega*mfbba; + //real Dxz = -c3o1*omega*mfbab; + //real Dyz = -c3o1*omega*mfabb; //3. // linear combinations diff --git a/src/gpu/VirtualFluids_GPU/GPU/Particles.cu b/src/gpu/VirtualFluids_GPU/GPU/Particles.cu index 193a1c2e7..f81d3f4e8 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/Particles.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/Particles.cu @@ -1929,370 +1929,373 @@ extern "C" __global__ void ParticleNoSlipDeviceComp27(real* coordX, unsigned int size_Mat, bool evenOrOdd) { - Distributions27 D; - if (evenOrOdd==true) - { - D.f[dirE ] = &DD[dirE *size_Mat]; - D.f[dirW ] = &DD[dirW *size_Mat]; - D.f[dirN ] = &DD[dirN *size_Mat]; - D.f[dirS ] = &DD[dirS *size_Mat]; - D.f[dirT ] = &DD[dirT *size_Mat]; - D.f[dirB ] = &DD[dirB *size_Mat]; - D.f[dirNE ] = &DD[dirNE *size_Mat]; - D.f[dirSW ] = &DD[dirSW *size_Mat]; - D.f[dirSE ] = &DD[dirSE *size_Mat]; - D.f[dirNW ] = &DD[dirNW *size_Mat]; - D.f[dirTE ] = &DD[dirTE *size_Mat]; - D.f[dirBW ] = &DD[dirBW *size_Mat]; - D.f[dirBE ] = &DD[dirBE *size_Mat]; - D.f[dirTW ] = &DD[dirTW *size_Mat]; - D.f[dirTN ] = &DD[dirTN *size_Mat]; - D.f[dirBS ] = &DD[dirBS *size_Mat]; - D.f[dirBN ] = &DD[dirBN *size_Mat]; - D.f[dirTS ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirTNE *size_Mat]; - D.f[dirTSW ] = &DD[dirTSW *size_Mat]; - D.f[dirTSE ] = &DD[dirTSE *size_Mat]; - D.f[dirTNW ] = &DD[dirTNW *size_Mat]; - D.f[dirBNE ] = &DD[dirBNE *size_Mat]; - D.f[dirBSW ] = &DD[dirBSW *size_Mat]; - D.f[dirBSE ] = &DD[dirBSE *size_Mat]; - D.f[dirBNW ] = &DD[dirBNW *size_Mat]; - } - else - { - D.f[dirW ] = &DD[dirE *size_Mat]; - D.f[dirE ] = &DD[dirW *size_Mat]; - D.f[dirS ] = &DD[dirN *size_Mat]; - D.f[dirN ] = &DD[dirS *size_Mat]; - D.f[dirB ] = &DD[dirT *size_Mat]; - D.f[dirT ] = &DD[dirB *size_Mat]; - D.f[dirSW ] = &DD[dirNE *size_Mat]; - D.f[dirNE ] = &DD[dirSW *size_Mat]; - D.f[dirNW ] = &DD[dirSE *size_Mat]; - D.f[dirSE ] = &DD[dirNW *size_Mat]; - D.f[dirBW ] = &DD[dirTE *size_Mat]; - D.f[dirTE ] = &DD[dirBW *size_Mat]; - D.f[dirTW ] = &DD[dirBE *size_Mat]; - D.f[dirBE ] = &DD[dirTW *size_Mat]; - D.f[dirBS ] = &DD[dirTN *size_Mat]; - D.f[dirTN ] = &DD[dirBS *size_Mat]; - D.f[dirTS ] = &DD[dirBN *size_Mat]; - D.f[dirBN ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirBSW *size_Mat]; - D.f[dirTSW ] = &DD[dirBNE *size_Mat]; - D.f[dirTSE ] = &DD[dirBNW *size_Mat]; - D.f[dirTNW ] = &DD[dirBSE *size_Mat]; - D.f[dirBNE ] = &DD[dirTSW *size_Mat]; - D.f[dirBSW ] = &DD[dirTNE *size_Mat]; - D.f[dirBSE ] = &DD[dirTNW *size_Mat]; - D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } - //////////////////////////////////////////////////////////////////////////////// - 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 nx = blockDim.x; - const unsigned ny = gridDim.x; - - const unsigned k = nx*(ny*z + y) + x; - ////////////////////////////////////////////////////////////////////////// - - if(k<sizeQ) - { - //////////////////////////////////////////////////////////////////////////////// - real *q_dirW, *q_dirS, *q_dirB; - // 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, - // *q_dirBE, *q_dirTW, *q_dirTN, *q_dirBS, *q_dirBN, *q_dirTS, - // *q_dirTNE, *q_dirTSW, *q_dirTSE, *q_dirTNW, *q_dirBNE, *q_dirBSW, - // *q_dirBSE, *q_dirBNW; - // q_dirE = &QQ[dirE *sizeQ]; - q_dirW = &QQ[dirW *sizeQ]; - // q_dirN = &QQ[dirN *sizeQ]; - q_dirS = &QQ[dirS *sizeQ]; - // q_dirT = &QQ[dirT *sizeQ]; - q_dirB = &QQ[dirB *sizeQ]; - // q_dirNE = &QQ[dirNE *sizeQ]; - // q_dirSW = &QQ[dirSW *sizeQ]; - // q_dirSE = &QQ[dirSE *sizeQ]; - // q_dirNW = &QQ[dirNW *sizeQ]; - // q_dirTE = &QQ[dirTE *sizeQ]; - // q_dirBW = &QQ[dirBW *sizeQ]; - // q_dirBE = &QQ[dirBE *sizeQ]; - // q_dirTW = &QQ[dirTW *sizeQ]; - // q_dirTN = &QQ[dirTN *sizeQ]; - // q_dirBS = &QQ[dirBS *sizeQ]; - // q_dirBN = &QQ[dirBN *sizeQ]; - // q_dirTS = &QQ[dirTS *sizeQ]; - // q_dirTNE = &QQ[dirTNE *sizeQ]; - // q_dirTSW = &QQ[dirTSW *sizeQ]; - // q_dirTSE = &QQ[dirTSE *sizeQ]; - // q_dirTNW = &QQ[dirTNW *sizeQ]; - // q_dirBNE = &QQ[dirBNE *sizeQ]; - // q_dirBSW = &QQ[dirBSW *sizeQ]; - // q_dirBSE = &QQ[dirBSE *sizeQ]; - // q_dirBNW = &QQ[dirBNW *sizeQ]; - //////////////////////////////////////////////////////////////////////////////// - // real *nx_dirE, *nx_dirW, *nx_dirN, *nx_dirS, *nx_dirT, *nx_dirB, - // *nx_dirNE, *nx_dirSW, *nx_dirSE, *nx_dirNW, *nx_dirTE, *nx_dirBW, - // *nx_dirBE, *nx_dirTW, *nx_dirTN, *nx_dirBS, *nx_dirBN, *nx_dirTS, - // *nx_dirTNE, *nx_dirTSW, *nx_dirTSE, *nx_dirTNW, *nx_dirBNE, *nx_dirBSW, - // *nx_dirBSE, *nx_dirBNW; - // nx_dirE = &NormalX[dirE *sizeQ]; - // nx_dirW = &NormalX[dirW *sizeQ]; - // nx_dirN = &NormalX[dirN *sizeQ]; - // nx_dirS = &NormalX[dirS *sizeQ]; - // nx_dirT = &NormalX[dirT *sizeQ]; - // nx_dirB = &NormalX[dirB *sizeQ]; - // nx_dirNE = &NormalX[dirNE *sizeQ]; - // nx_dirSW = &NormalX[dirSW *sizeQ]; - // nx_dirSE = &NormalX[dirSE *sizeQ]; - // nx_dirNW = &NormalX[dirNW *sizeQ]; - // nx_dirTE = &NormalX[dirTE *sizeQ]; - // nx_dirBW = &NormalX[dirBW *sizeQ]; - // nx_dirBE = &NormalX[dirBE *sizeQ]; - // nx_dirTW = &NormalX[dirTW *sizeQ]; - // nx_dirTN = &NormalX[dirTN *sizeQ]; - // nx_dirBS = &NormalX[dirBS *sizeQ]; - // nx_dirBN = &NormalX[dirBN *sizeQ]; - // nx_dirTS = &NormalX[dirTS *sizeQ]; - // nx_dirTNE = &NormalX[dirTNE *sizeQ]; - // nx_dirTSW = &NormalX[dirTSW *sizeQ]; - // nx_dirTSE = &NormalX[dirTSE *sizeQ]; - // nx_dirTNW = &NormalX[dirTNW *sizeQ]; - // nx_dirBNE = &NormalX[dirBNE *sizeQ]; - // nx_dirBSW = &NormalX[dirBSW *sizeQ]; - // nx_dirBSE = &NormalX[dirBSE *sizeQ]; - // nx_dirBNW = &NormalX[dirBNW *sizeQ]; - //////////////////////////////////////////////////////////////////////////////// - // real *ny_dirE, *ny_dirW, *ny_dirN, *ny_dirS, *ny_dirT, *ny_dirB, - // *ny_dirNE, *ny_dirSW, *ny_dirSE, *ny_dirNW, *ny_dirTE, *ny_dirBW, - // *ny_dirBE, *ny_dirTW, *ny_dirTN, *ny_dirBS, *ny_dirBN, *ny_dirTS, - // *ny_dirTNE, *ny_dirTSW, *ny_dirTSE, *ny_dirTNW, *ny_dirBNE, *ny_dirBSW, - // *ny_dirBSE, *ny_dirBNW; - // ny_dirE = &NormalY[dirE *sizeQ]; - // ny_dirW = &NormalY[dirW *sizeQ]; - // ny_dirN = &NormalY[dirN *sizeQ]; - // ny_dirS = &NormalY[dirS *sizeQ]; - // ny_dirT = &NormalY[dirT *sizeQ]; - // ny_dirB = &NormalY[dirB *sizeQ]; - // ny_dirNE = &NormalY[dirNE *sizeQ]; - // ny_dirSW = &NormalY[dirSW *sizeQ]; - // ny_dirSE = &NormalY[dirSE *sizeQ]; - // ny_dirNW = &NormalY[dirNW *sizeQ]; - // ny_dirTE = &NormalY[dirTE *sizeQ]; - // ny_dirBW = &NormalY[dirBW *sizeQ]; - // ny_dirBE = &NormalY[dirBE *sizeQ]; - // ny_dirTW = &NormalY[dirTW *sizeQ]; - // ny_dirTN = &NormalY[dirTN *sizeQ]; - // ny_dirBS = &NormalY[dirBS *sizeQ]; - // ny_dirBN = &NormalY[dirBN *sizeQ]; - // ny_dirTS = &NormalY[dirTS *sizeQ]; - // ny_dirTNE = &NormalY[dirTNE *sizeQ]; - // ny_dirTSW = &NormalY[dirTSW *sizeQ]; - // ny_dirTSE = &NormalY[dirTSE *sizeQ]; - // ny_dirTNW = &NormalY[dirTNW *sizeQ]; - // ny_dirBNE = &NormalY[dirBNE *sizeQ]; - // ny_dirBSW = &NormalY[dirBSW *sizeQ]; - // ny_dirBSE = &NormalY[dirBSE *sizeQ]; - // ny_dirBNW = &NormalY[dirBNW *sizeQ]; - //////////////////////////////////////////////////////////////////////////////// - // real *nz_dirE, *nz_dirW, *nz_dirN, *nz_dirS, *nz_dirT, *nz_dirB, - // *nz_dirNE, *nz_dirSW, *nz_dirSE, *nz_dirNW, *nz_dirTE, *nz_dirBW, - // *nz_dirBE, *nz_dirTW, *nz_dirTN, *nz_dirBS, *nz_dirBN, *nz_dirTS, - // *nz_dirTNE, *nz_dirTSW, *nz_dirTSE, *nz_dirTNW, *nz_dirBNE, *nz_dirBSW, - // *nz_dirBSE, *nz_dirBNW; - // nz_dirE = &NormalZ[dirE *sizeQ]; - // nz_dirW = &NormalZ[dirW *sizeQ]; - // nz_dirN = &NormalZ[dirN *sizeQ]; - // nz_dirS = &NormalZ[dirS *sizeQ]; - // nz_dirT = &NormalZ[dirT *sizeQ]; - // nz_dirB = &NormalZ[dirB *sizeQ]; - // nz_dirNE = &NormalZ[dirNE *sizeQ]; - // nz_dirSW = &NormalZ[dirSW *sizeQ]; - // nz_dirSE = &NormalZ[dirSE *sizeQ]; - // nz_dirNW = &NormalZ[dirNW *sizeQ]; - // nz_dirTE = &NormalZ[dirTE *sizeQ]; - // nz_dirBW = &NormalZ[dirBW *sizeQ]; - // nz_dirBE = &NormalZ[dirBE *sizeQ]; - // nz_dirTW = &NormalZ[dirTW *sizeQ]; - // nz_dirTN = &NormalZ[dirTN *sizeQ]; - // nz_dirBS = &NormalZ[dirBS *sizeQ]; - // nz_dirBN = &NormalZ[dirBN *sizeQ]; - // nz_dirTS = &NormalZ[dirTS *sizeQ]; - // nz_dirTNE = &NormalZ[dirTNE *sizeQ]; - // nz_dirTSW = &NormalZ[dirTSW *sizeQ]; - // nz_dirTSE = &NormalZ[dirTSE *sizeQ]; - // nz_dirTNW = &NormalZ[dirTNW *sizeQ]; - // nz_dirBNE = &NormalZ[dirBNE *sizeQ]; - // nz_dirBSW = &NormalZ[dirBSW *sizeQ]; - // nz_dirBSE = &NormalZ[dirBSE *sizeQ]; - // nz_dirBNW = &NormalZ[dirBNW *sizeQ]; - //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - bool changeCell = false; - unsigned int KQK = k_Q[k]; - //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - if( q_dirW[k] > c0o1 || q_dirS[k] > c0o1 || q_dirB[k] > c0o1 ) { - KQK = neighborWSB[KQK]; - changeCell = true; - } - if( q_dirW[k] == c0o1 && changeCell == true ) { - KQK = neighborX[KQK]; - } - if( q_dirS[k] == c0o1 && changeCell == true ) { - KQK = neighborY[KQK]; - } - if( q_dirB[k] == c0o1 && changeCell == true ) { - KQK = neighborZ[KQK]; - } - - for(int i = 0; i < numberOfParticles; i++){ - //push back? - } - - //////////////////////////////////////////////////////////////////////////////// - //index - //unsigned int KQK = k_Q[k]; - unsigned int kzero= KQK; - unsigned int ke = KQK; - unsigned int kw = neighborX[KQK]; - unsigned int kn = KQK; - unsigned int ks = neighborY[KQK]; - unsigned int kt = KQK; - unsigned int kb = neighborZ[KQK]; - unsigned int ksw = neighborY[kw]; - unsigned int kne = KQK; - unsigned int kse = ks; - unsigned int knw = kw; - unsigned int kbw = neighborZ[kw]; - unsigned int kte = KQK; - unsigned int kbe = kb; - unsigned int ktw = kw; - unsigned int kbs = neighborZ[ks]; - unsigned int ktn = KQK; - unsigned int kbn = kb; - unsigned int kts = ks; - unsigned int ktse = ks; - unsigned int kbnw = kbw; - unsigned int ktnw = kw; - unsigned int kbse = kbs; - unsigned int ktsw = ksw; - unsigned int kbne = kb; - unsigned int ktne = KQK; - unsigned int kbsw = neighborZ[ksw]; - //////////////////////////////////////////////////////////////////////////////// - real f_W = (D.f[dirE ])[ke ]; - real f_E = (D.f[dirW ])[kw ]; - real f_S = (D.f[dirN ])[kn ]; - real f_N = (D.f[dirS ])[ks ]; - real f_B = (D.f[dirT ])[kt ]; - real f_T = (D.f[dirB ])[kb ]; - real f_SW = (D.f[dirNE ])[kne ]; - real f_NE = (D.f[dirSW ])[ksw ]; - real f_NW = (D.f[dirSE ])[kse ]; - real f_SE = (D.f[dirNW ])[knw ]; - real f_BW = (D.f[dirTE ])[kte ]; - real f_TE = (D.f[dirBW ])[kbw ]; - real f_TW = (D.f[dirBE ])[kbe ]; - real f_BE = (D.f[dirTW ])[ktw ]; - real f_BS = (D.f[dirTN ])[ktn ]; - real f_TN = (D.f[dirBS ])[kbs ]; - real f_TS = (D.f[dirBN ])[kbn ]; - real f_BN = (D.f[dirTS ])[kts ]; - real f_BSW = (D.f[dirTNE ])[ktne ]; - real f_BNE = (D.f[dirTSW ])[ktsw ]; - real f_BNW = (D.f[dirTSE ])[ktse ]; - real f_BSE = (D.f[dirTNW ])[ktnw ]; - real f_TSW = (D.f[dirBNE ])[kbne ]; - real f_TNE = (D.f[dirBSW ])[kbsw ]; - real f_TNW = (D.f[dirBSE ])[kbse ]; - real f_TSE = (D.f[dirBNW ])[kbnw ]; - //////////////////////////////////////////////////////////////////////////////// - // real feq, q; - real vx1, vx2, vx3, drho; - drho = f_TSE + f_TNW + f_TNE + f_TSW + f_BSE + f_BNW + f_BNE + f_BSW + - f_BN + f_TS + f_TN + f_BS + f_BE + f_TW + f_TE + f_BW + f_SE + f_NW + f_NE + f_SW + - f_T + f_B + f_N + f_S + f_E + f_W + ((D.f[dirZERO])[kzero]); - - vx1 = (((f_TSE - f_BNW) - (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + - ((f_BE - f_TW) + (f_TE - f_BW)) + ((f_SE - f_NW) + (f_NE - f_SW)) + - (f_E - f_W)) / (c1o1 + drho); - - - vx2 = ((-(f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + - ((f_BN - f_TS) + (f_TN - f_BS)) + (-(f_SE - f_NW) + (f_NE - f_SW)) + - (f_N - f_S)) / (c1o1 + drho); - - vx3 = (((f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) + (f_TSW - f_BNE)) + - (-(f_BN - f_TS) + (f_TN - f_BS)) + ((f_TE - f_BW) - (f_BE - f_TW)) + - (f_T - f_B)) / (c1o1 + drho); - - real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3) * (c1o1 + drho); - - ////////////////////////////////////////////////////////////////////////// - if (evenOrOdd==false) - { - D.f[dirE ] = &DD[dirE *size_Mat]; - D.f[dirW ] = &DD[dirW *size_Mat]; - D.f[dirN ] = &DD[dirN *size_Mat]; - D.f[dirS ] = &DD[dirS *size_Mat]; - D.f[dirT ] = &DD[dirT *size_Mat]; - D.f[dirB ] = &DD[dirB *size_Mat]; - D.f[dirNE ] = &DD[dirNE *size_Mat]; - D.f[dirSW ] = &DD[dirSW *size_Mat]; - D.f[dirSE ] = &DD[dirSE *size_Mat]; - D.f[dirNW ] = &DD[dirNW *size_Mat]; - D.f[dirTE ] = &DD[dirTE *size_Mat]; - D.f[dirBW ] = &DD[dirBW *size_Mat]; - D.f[dirBE ] = &DD[dirBE *size_Mat]; - D.f[dirTW ] = &DD[dirTW *size_Mat]; - D.f[dirTN ] = &DD[dirTN *size_Mat]; - D.f[dirBS ] = &DD[dirBS *size_Mat]; - D.f[dirBN ] = &DD[dirBN *size_Mat]; - D.f[dirTS ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirTNE *size_Mat]; - D.f[dirTSW ] = &DD[dirTSW *size_Mat]; - D.f[dirTSE ] = &DD[dirTSE *size_Mat]; - D.f[dirTNW ] = &DD[dirTNW *size_Mat]; - D.f[dirBNE ] = &DD[dirBNE *size_Mat]; - D.f[dirBSW ] = &DD[dirBSW *size_Mat]; - D.f[dirBSE ] = &DD[dirBSE *size_Mat]; - D.f[dirBNW ] = &DD[dirBNW *size_Mat]; - } - else - { - D.f[dirW ] = &DD[dirE *size_Mat]; - D.f[dirE ] = &DD[dirW *size_Mat]; - D.f[dirS ] = &DD[dirN *size_Mat]; - D.f[dirN ] = &DD[dirS *size_Mat]; - D.f[dirB ] = &DD[dirT *size_Mat]; - D.f[dirT ] = &DD[dirB *size_Mat]; - D.f[dirSW ] = &DD[dirNE *size_Mat]; - D.f[dirNE ] = &DD[dirSW *size_Mat]; - D.f[dirNW ] = &DD[dirSE *size_Mat]; - D.f[dirSE ] = &DD[dirNW *size_Mat]; - D.f[dirBW ] = &DD[dirTE *size_Mat]; - D.f[dirTE ] = &DD[dirBW *size_Mat]; - D.f[dirTW ] = &DD[dirBE *size_Mat]; - D.f[dirBE ] = &DD[dirTW *size_Mat]; - D.f[dirBS ] = &DD[dirTN *size_Mat]; - D.f[dirTN ] = &DD[dirBS *size_Mat]; - D.f[dirTS ] = &DD[dirBN *size_Mat]; - D.f[dirBN ] = &DD[dirTS *size_Mat]; - D.f[dirZERO] = &DD[dirZERO*size_Mat]; - D.f[dirTNE ] = &DD[dirBSW *size_Mat]; - D.f[dirTSW ] = &DD[dirBNE *size_Mat]; - D.f[dirTSE ] = &DD[dirBNW *size_Mat]; - D.f[dirTNW ] = &DD[dirBSE *size_Mat]; - D.f[dirBNE ] = &DD[dirTSW *size_Mat]; - D.f[dirBSW ] = &DD[dirTNE *size_Mat]; - D.f[dirBSE ] = &DD[dirTNW *size_Mat]; - D.f[dirBNW ] = &DD[dirTSE *size_Mat]; - } - } + //TODO: What is this function for??? + + //Distributions27 D; + //if (evenOrOdd==true) + //{ + // D.f[dirE ] = &DD[dirE *size_Mat]; + // D.f[dirW ] = &DD[dirW *size_Mat]; + // D.f[dirN ] = &DD[dirN *size_Mat]; + // D.f[dirS ] = &DD[dirS *size_Mat]; + // D.f[dirT ] = &DD[dirT *size_Mat]; + // D.f[dirB ] = &DD[dirB *size_Mat]; + // D.f[dirNE ] = &DD[dirNE *size_Mat]; + // D.f[dirSW ] = &DD[dirSW *size_Mat]; + // D.f[dirSE ] = &DD[dirSE *size_Mat]; + // D.f[dirNW ] = &DD[dirNW *size_Mat]; + // D.f[dirTE ] = &DD[dirTE *size_Mat]; + // D.f[dirBW ] = &DD[dirBW *size_Mat]; + // D.f[dirBE ] = &DD[dirBE *size_Mat]; + // D.f[dirTW ] = &DD[dirTW *size_Mat]; + // D.f[dirTN ] = &DD[dirTN *size_Mat]; + // D.f[dirBS ] = &DD[dirBS *size_Mat]; + // D.f[dirBN ] = &DD[dirBN *size_Mat]; + // D.f[dirTS ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirTNE *size_Mat]; + // D.f[dirTSW ] = &DD[dirTSW *size_Mat]; + // D.f[dirTSE ] = &DD[dirTSE *size_Mat]; + // D.f[dirTNW ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNE ] = &DD[dirBNE *size_Mat]; + // D.f[dirBSW ] = &DD[dirBSW *size_Mat]; + // D.f[dirBSE ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNW ] = &DD[dirBNW *size_Mat]; + //} + //else + //{ + // D.f[dirW ] = &DD[dirE *size_Mat]; + // D.f[dirE ] = &DD[dirW *size_Mat]; + // D.f[dirS ] = &DD[dirN *size_Mat]; + // D.f[dirN ] = &DD[dirS *size_Mat]; + // D.f[dirB ] = &DD[dirT *size_Mat]; + // D.f[dirT ] = &DD[dirB *size_Mat]; + // D.f[dirSW ] = &DD[dirNE *size_Mat]; + // D.f[dirNE ] = &DD[dirSW *size_Mat]; + // D.f[dirNW ] = &DD[dirSE *size_Mat]; + // D.f[dirSE ] = &DD[dirNW *size_Mat]; + // D.f[dirBW ] = &DD[dirTE *size_Mat]; + // D.f[dirTE ] = &DD[dirBW *size_Mat]; + // D.f[dirTW ] = &DD[dirBE *size_Mat]; + // D.f[dirBE ] = &DD[dirTW *size_Mat]; + // D.f[dirBS ] = &DD[dirTN *size_Mat]; + // D.f[dirTN ] = &DD[dirBS *size_Mat]; + // D.f[dirTS ] = &DD[dirBN *size_Mat]; + // D.f[dirBN ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirBSW *size_Mat]; + // D.f[dirTSW ] = &DD[dirBNE *size_Mat]; + // D.f[dirTSE ] = &DD[dirBNW *size_Mat]; + // D.f[dirTNW ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNE ] = &DD[dirTSW *size_Mat]; + // D.f[dirBSW ] = &DD[dirTNE *size_Mat]; + // D.f[dirBSE ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNW ] = &DD[dirTSE *size_Mat]; + //} + ////////////////////////////////////////////////////////////////////////////////// + //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 nx = blockDim.x; + //const unsigned ny = gridDim.x; + + //const unsigned k = nx*(ny*z + y) + x; + //////////////////////////////////////////////////////////////////////////// + + //if(k < sizeQ) + //{ + // //////////////////////////////////////////////////////////////////////////////// + // real *q_dirW, *q_dirS, *q_dirB; + // // 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, + // // *q_dirBE, *q_dirTW, *q_dirTN, *q_dirBS, *q_dirBN, *q_dirTS, + // // *q_dirTNE, *q_dirTSW, *q_dirTSE, *q_dirTNW, *q_dirBNE, *q_dirBSW, + // // *q_dirBSE, *q_dirBNW; + // // q_dirE = &QQ[dirE *sizeQ]; + // q_dirW = &QQ[dirW *sizeQ]; + // // q_dirN = &QQ[dirN *sizeQ]; + // q_dirS = &QQ[dirS *sizeQ]; + // // q_dirT = &QQ[dirT *sizeQ]; + // q_dirB = &QQ[dirB *sizeQ]; + // // q_dirNE = &QQ[dirNE *sizeQ]; + // // q_dirSW = &QQ[dirSW *sizeQ]; + // // q_dirSE = &QQ[dirSE *sizeQ]; + // // q_dirNW = &QQ[dirNW *sizeQ]; + // // q_dirTE = &QQ[dirTE *sizeQ]; + // // q_dirBW = &QQ[dirBW *sizeQ]; + // // q_dirBE = &QQ[dirBE *sizeQ]; + // // q_dirTW = &QQ[dirTW *sizeQ]; + // // q_dirTN = &QQ[dirTN *sizeQ]; + // // q_dirBS = &QQ[dirBS *sizeQ]; + // // q_dirBN = &QQ[dirBN *sizeQ]; + // // q_dirTS = &QQ[dirTS *sizeQ]; + // // q_dirTNE = &QQ[dirTNE *sizeQ]; + // // q_dirTSW = &QQ[dirTSW *sizeQ]; + // // q_dirTSE = &QQ[dirTSE *sizeQ]; + // // q_dirTNW = &QQ[dirTNW *sizeQ]; + // // q_dirBNE = &QQ[dirBNE *sizeQ]; + // // q_dirBSW = &QQ[dirBSW *sizeQ]; + // // q_dirBSE = &QQ[dirBSE *sizeQ]; + // // q_dirBNW = &QQ[dirBNW *sizeQ]; + // //////////////////////////////////////////////////////////////////////////////// + // // real *nx_dirE, *nx_dirW, *nx_dirN, *nx_dirS, *nx_dirT, *nx_dirB, + // // *nx_dirNE, *nx_dirSW, *nx_dirSE, *nx_dirNW, *nx_dirTE, *nx_dirBW, + // // *nx_dirBE, *nx_dirTW, *nx_dirTN, *nx_dirBS, *nx_dirBN, *nx_dirTS, + // // *nx_dirTNE, *nx_dirTSW, *nx_dirTSE, *nx_dirTNW, *nx_dirBNE, *nx_dirBSW, + // // *nx_dirBSE, *nx_dirBNW; + // // nx_dirE = &NormalX[dirE *sizeQ]; + // // nx_dirW = &NormalX[dirW *sizeQ]; + // // nx_dirN = &NormalX[dirN *sizeQ]; + // // nx_dirS = &NormalX[dirS *sizeQ]; + // // nx_dirT = &NormalX[dirT *sizeQ]; + // // nx_dirB = &NormalX[dirB *sizeQ]; + // // nx_dirNE = &NormalX[dirNE *sizeQ]; + // // nx_dirSW = &NormalX[dirSW *sizeQ]; + // // nx_dirSE = &NormalX[dirSE *sizeQ]; + // // nx_dirNW = &NormalX[dirNW *sizeQ]; + // // nx_dirTE = &NormalX[dirTE *sizeQ]; + // // nx_dirBW = &NormalX[dirBW *sizeQ]; + // // nx_dirBE = &NormalX[dirBE *sizeQ]; + // // nx_dirTW = &NormalX[dirTW *sizeQ]; + // // nx_dirTN = &NormalX[dirTN *sizeQ]; + // // nx_dirBS = &NormalX[dirBS *sizeQ]; + // // nx_dirBN = &NormalX[dirBN *sizeQ]; + // // nx_dirTS = &NormalX[dirTS *sizeQ]; + // // nx_dirTNE = &NormalX[dirTNE *sizeQ]; + // // nx_dirTSW = &NormalX[dirTSW *sizeQ]; + // // nx_dirTSE = &NormalX[dirTSE *sizeQ]; + // // nx_dirTNW = &NormalX[dirTNW *sizeQ]; + // // nx_dirBNE = &NormalX[dirBNE *sizeQ]; + // // nx_dirBSW = &NormalX[dirBSW *sizeQ]; + // // nx_dirBSE = &NormalX[dirBSE *sizeQ]; + // // nx_dirBNW = &NormalX[dirBNW *sizeQ]; + // //////////////////////////////////////////////////////////////////////////////// + // // real *ny_dirE, *ny_dirW, *ny_dirN, *ny_dirS, *ny_dirT, *ny_dirB, + // // *ny_dirNE, *ny_dirSW, *ny_dirSE, *ny_dirNW, *ny_dirTE, *ny_dirBW, + // // *ny_dirBE, *ny_dirTW, *ny_dirTN, *ny_dirBS, *ny_dirBN, *ny_dirTS, + // // *ny_dirTNE, *ny_dirTSW, *ny_dirTSE, *ny_dirTNW, *ny_dirBNE, *ny_dirBSW, + // // *ny_dirBSE, *ny_dirBNW; + // // ny_dirE = &NormalY[dirE *sizeQ]; + // // ny_dirW = &NormalY[dirW *sizeQ]; + // // ny_dirN = &NormalY[dirN *sizeQ]; + // // ny_dirS = &NormalY[dirS *sizeQ]; + // // ny_dirT = &NormalY[dirT *sizeQ]; + // // ny_dirB = &NormalY[dirB *sizeQ]; + // // ny_dirNE = &NormalY[dirNE *sizeQ]; + // // ny_dirSW = &NormalY[dirSW *sizeQ]; + // // ny_dirSE = &NormalY[dirSE *sizeQ]; + // // ny_dirNW = &NormalY[dirNW *sizeQ]; + // // ny_dirTE = &NormalY[dirTE *sizeQ]; + // // ny_dirBW = &NormalY[dirBW *sizeQ]; + // // ny_dirBE = &NormalY[dirBE *sizeQ]; + // // ny_dirTW = &NormalY[dirTW *sizeQ]; + // // ny_dirTN = &NormalY[dirTN *sizeQ]; + // // ny_dirBS = &NormalY[dirBS *sizeQ]; + // // ny_dirBN = &NormalY[dirBN *sizeQ]; + // // ny_dirTS = &NormalY[dirTS *sizeQ]; + // // ny_dirTNE = &NormalY[dirTNE *sizeQ]; + // // ny_dirTSW = &NormalY[dirTSW *sizeQ]; + // // ny_dirTSE = &NormalY[dirTSE *sizeQ]; + // // ny_dirTNW = &NormalY[dirTNW *sizeQ]; + // // ny_dirBNE = &NormalY[dirBNE *sizeQ]; + // // ny_dirBSW = &NormalY[dirBSW *sizeQ]; + // // ny_dirBSE = &NormalY[dirBSE *sizeQ]; + // // ny_dirBNW = &NormalY[dirBNW *sizeQ]; + // //////////////////////////////////////////////////////////////////////////////// + // // real *nz_dirE, *nz_dirW, *nz_dirN, *nz_dirS, *nz_dirT, *nz_dirB, + // // *nz_dirNE, *nz_dirSW, *nz_dirSE, *nz_dirNW, *nz_dirTE, *nz_dirBW, + // // *nz_dirBE, *nz_dirTW, *nz_dirTN, *nz_dirBS, *nz_dirBN, *nz_dirTS, + // // *nz_dirTNE, *nz_dirTSW, *nz_dirTSE, *nz_dirTNW, *nz_dirBNE, *nz_dirBSW, + // // *nz_dirBSE, *nz_dirBNW; + // // nz_dirE = &NormalZ[dirE *sizeQ]; + // // nz_dirW = &NormalZ[dirW *sizeQ]; + // // nz_dirN = &NormalZ[dirN *sizeQ]; + // // nz_dirS = &NormalZ[dirS *sizeQ]; + // // nz_dirT = &NormalZ[dirT *sizeQ]; + // // nz_dirB = &NormalZ[dirB *sizeQ]; + // // nz_dirNE = &NormalZ[dirNE *sizeQ]; + // // nz_dirSW = &NormalZ[dirSW *sizeQ]; + // // nz_dirSE = &NormalZ[dirSE *sizeQ]; + // // nz_dirNW = &NormalZ[dirNW *sizeQ]; + // // nz_dirTE = &NormalZ[dirTE *sizeQ]; + // // nz_dirBW = &NormalZ[dirBW *sizeQ]; + // // nz_dirBE = &NormalZ[dirBE *sizeQ]; + // // nz_dirTW = &NormalZ[dirTW *sizeQ]; + // // nz_dirTN = &NormalZ[dirTN *sizeQ]; + // // nz_dirBS = &NormalZ[dirBS *sizeQ]; + // // nz_dirBN = &NormalZ[dirBN *sizeQ]; + // // nz_dirTS = &NormalZ[dirTS *sizeQ]; + // // nz_dirTNE = &NormalZ[dirTNE *sizeQ]; + // // nz_dirTSW = &NormalZ[dirTSW *sizeQ]; + // // nz_dirTSE = &NormalZ[dirTSE *sizeQ]; + // // nz_dirTNW = &NormalZ[dirTNW *sizeQ]; + // // nz_dirBNE = &NormalZ[dirBNE *sizeQ]; + // // nz_dirBSW = &NormalZ[dirBSW *sizeQ]; + // // nz_dirBSE = &NormalZ[dirBSE *sizeQ]; + // // nz_dirBNW = &NormalZ[dirBNW *sizeQ]; + // //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// + //bool changeCell = false; + // unsigned int KQK = k_Q[k]; + // //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// + //if( q_dirW[k] > c0o1 || q_dirS[k] > c0o1 || q_dirB[k] > c0o1 ) { + // KQK = neighborWSB[KQK]; + // changeCell = true; + //} + //if( q_dirW[k] == c0o1 && changeCell == true ) { + // KQK = neighborX[KQK]; + //} + //if( q_dirS[k] == c0o1 && changeCell == true ) { + // KQK = neighborY[KQK]; + //} + //if( q_dirB[k] == c0o1 && changeCell == true ) { + // KQK = neighborZ[KQK]; + //} + + ////for(int i = 0; i < numberOfParticles; i++){ + //// //push back? + ////} + + // //////////////////////////////////////////////////////////////////////////////// + // //index + // //unsigned int KQK = k_Q[k]; + // unsigned int kzero= KQK; + // unsigned int ke = KQK; + // unsigned int kw = neighborX[KQK]; + // unsigned int kn = KQK; + // unsigned int ks = neighborY[KQK]; + // unsigned int kt = KQK; + // unsigned int kb = neighborZ[KQK]; + // unsigned int ksw = neighborY[kw]; + // unsigned int kne = KQK; + // unsigned int kse = ks; + // unsigned int knw = kw; + // unsigned int kbw = neighborZ[kw]; + // unsigned int kte = KQK; + // unsigned int kbe = kb; + // unsigned int ktw = kw; + // unsigned int kbs = neighborZ[ks]; + // unsigned int ktn = KQK; + // unsigned int kbn = kb; + // unsigned int kts = ks; + // unsigned int ktse = ks; + // unsigned int kbnw = kbw; + // unsigned int ktnw = kw; + // unsigned int kbse = kbs; + // unsigned int ktsw = ksw; + // unsigned int kbne = kb; + // unsigned int ktne = KQK; + // unsigned int kbsw = neighborZ[ksw]; + // //////////////////////////////////////////////////////////////////////////////// + // real f_W = (D.f[dirE ])[ke ]; + // real f_E = (D.f[dirW ])[kw ]; + // real f_S = (D.f[dirN ])[kn ]; + // real f_N = (D.f[dirS ])[ks ]; + // real f_B = (D.f[dirT ])[kt ]; + // real f_T = (D.f[dirB ])[kb ]; + // real f_SW = (D.f[dirNE ])[kne ]; + // real f_NE = (D.f[dirSW ])[ksw ]; + // real f_NW = (D.f[dirSE ])[kse ]; + // real f_SE = (D.f[dirNW ])[knw ]; + // real f_BW = (D.f[dirTE ])[kte ]; + // real f_TE = (D.f[dirBW ])[kbw ]; + // real f_TW = (D.f[dirBE ])[kbe ]; + // real f_BE = (D.f[dirTW ])[ktw ]; + // real f_BS = (D.f[dirTN ])[ktn ]; + // real f_TN = (D.f[dirBS ])[kbs ]; + // real f_TS = (D.f[dirBN ])[kbn ]; + // real f_BN = (D.f[dirTS ])[kts ]; + // real f_BSW = (D.f[dirTNE ])[ktne ]; + // real f_BNE = (D.f[dirTSW ])[ktsw ]; + // real f_BNW = (D.f[dirTSE ])[ktse ]; + // real f_BSE = (D.f[dirTNW ])[ktnw ]; + // real f_TSW = (D.f[dirBNE ])[kbne ]; + // real f_TNE = (D.f[dirBSW ])[kbsw ]; + // real f_TNW = (D.f[dirBSE ])[kbse ]; + // real f_TSE = (D.f[dirBNW ])[kbnw ]; + // //////////////////////////////////////////////////////////////////////////////// + // // real feq, q; + // real vx1, vx2, vx3, drho; + // drho = f_TSE + f_TNW + f_TNE + f_TSW + f_BSE + f_BNW + f_BNE + f_BSW + + // f_BN + f_TS + f_TN + f_BS + f_BE + f_TW + f_TE + f_BW + f_SE + f_NW + f_NE + f_SW + + // f_T + f_B + f_N + f_S + f_E + f_W + ((D.f[dirZERO])[kzero]); + + // vx1 = (((f_TSE - f_BNW) - (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + + // ((f_BE - f_TW) + (f_TE - f_BW)) + ((f_SE - f_NW) + (f_NE - f_SW)) + + // (f_E - f_W)) / (c1o1 + drho); + // + + // vx2 = ((-(f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) - (f_TSW - f_BNE)) + + // ((f_BN - f_TS) + (f_TN - f_BS)) + (-(f_SE - f_NW) + (f_NE - f_SW)) + + // (f_N - f_S)) / (c1o1 + drho); + + // vx3 = (((f_TSE - f_BNW) + (f_TNW - f_BSE)) + ((f_TNE - f_BSW) + (f_TSW - f_BNE)) + + // (-(f_BN - f_TS) + (f_TN - f_BS)) + ((f_TE - f_BW) - (f_BE - f_TW)) + + // (f_T - f_B)) / (c1o1 + drho); + + // //real cu_sq=c3o2*(vx1*vx1+vx2*vx2+vx3*vx3) * (c1o1 + drho); + + // ////////////////////////////////////////////////////////////////////////// + // if (evenOrOdd==false) + // { + // D.f[dirE ] = &DD[dirE *size_Mat]; + // D.f[dirW ] = &DD[dirW *size_Mat]; + // D.f[dirN ] = &DD[dirN *size_Mat]; + // D.f[dirS ] = &DD[dirS *size_Mat]; + // D.f[dirT ] = &DD[dirT *size_Mat]; + // D.f[dirB ] = &DD[dirB *size_Mat]; + // D.f[dirNE ] = &DD[dirNE *size_Mat]; + // D.f[dirSW ] = &DD[dirSW *size_Mat]; + // D.f[dirSE ] = &DD[dirSE *size_Mat]; + // D.f[dirNW ] = &DD[dirNW *size_Mat]; + // D.f[dirTE ] = &DD[dirTE *size_Mat]; + // D.f[dirBW ] = &DD[dirBW *size_Mat]; + // D.f[dirBE ] = &DD[dirBE *size_Mat]; + // D.f[dirTW ] = &DD[dirTW *size_Mat]; + // D.f[dirTN ] = &DD[dirTN *size_Mat]; + // D.f[dirBS ] = &DD[dirBS *size_Mat]; + // D.f[dirBN ] = &DD[dirBN *size_Mat]; + // D.f[dirTS ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirTNE *size_Mat]; + // D.f[dirTSW ] = &DD[dirTSW *size_Mat]; + // D.f[dirTSE ] = &DD[dirTSE *size_Mat]; + // D.f[dirTNW ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNE ] = &DD[dirBNE *size_Mat]; + // D.f[dirBSW ] = &DD[dirBSW *size_Mat]; + // D.f[dirBSE ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNW ] = &DD[dirBNW *size_Mat]; + // } + // else + // { + // D.f[dirW ] = &DD[dirE *size_Mat]; + // D.f[dirE ] = &DD[dirW *size_Mat]; + // D.f[dirS ] = &DD[dirN *size_Mat]; + // D.f[dirN ] = &DD[dirS *size_Mat]; + // D.f[dirB ] = &DD[dirT *size_Mat]; + // D.f[dirT ] = &DD[dirB *size_Mat]; + // D.f[dirSW ] = &DD[dirNE *size_Mat]; + // D.f[dirNE ] = &DD[dirSW *size_Mat]; + // D.f[dirNW ] = &DD[dirSE *size_Mat]; + // D.f[dirSE ] = &DD[dirNW *size_Mat]; + // D.f[dirBW ] = &DD[dirTE *size_Mat]; + // D.f[dirTE ] = &DD[dirBW *size_Mat]; + // D.f[dirTW ] = &DD[dirBE *size_Mat]; + // D.f[dirBE ] = &DD[dirTW *size_Mat]; + // D.f[dirBS ] = &DD[dirTN *size_Mat]; + // D.f[dirTN ] = &DD[dirBS *size_Mat]; + // D.f[dirTS ] = &DD[dirBN *size_Mat]; + // D.f[dirBN ] = &DD[dirTS *size_Mat]; + // D.f[dirZERO] = &DD[dirZERO*size_Mat]; + // D.f[dirTNE ] = &DD[dirBSW *size_Mat]; + // D.f[dirTSW ] = &DD[dirBNE *size_Mat]; + // D.f[dirTSE ] = &DD[dirBNW *size_Mat]; + // D.f[dirTNW ] = &DD[dirBSE *size_Mat]; + // D.f[dirBNE ] = &DD[dirTSW *size_Mat]; + // D.f[dirBSW ] = &DD[dirTNE *size_Mat]; + // D.f[dirBSE ] = &DD[dirTNW *size_Mat]; + // D.f[dirBNW ] = &DD[dirTSE *size_Mat]; + // } + //} } diff --git a/src/gpu/VirtualFluids_GPU/GPU/ScaleFC27.cu b/src/gpu/VirtualFluids_GPU/GPU/ScaleFC27.cu index 90b778ec2..4b085a7d6 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/ScaleFC27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/ScaleFC27.cu @@ -772,7 +772,7 @@ extern "C" __global__ void scaleFC_0817_comp_27( real* DC, real m0, m1, m2, oMdrho; real mxxPyyPzz, mxxMyy, mxxMzz, mxxyPyzz, mxxyMyzz, mxxzPyyz, mxxzMyyz, mxyyPxzz, mxyyMxzz; //real qudricLimit = c1o100;//ganz schlechte Idee -> muss global sein - real O3 = c2o1 - o; + //real O3 = c2o1 - o; //real residu, residutmp; //residutmp = c0o1;///*-*/ c2o9 * (1./o - c1o2) * eps_new * eps_new; real NeqOn = c1o1;//zero;//one; //.... one = on ..... zero = off @@ -10504,7 +10504,7 @@ extern "C" __global__ void scaleFC_RhoSq_comp_27(real* DC, real m0, m1, m2, vvx, vvy, vvz, vx2, vy2, vz2, oMdrho; real mxxPyyPzz, mxxMyy, mxxMzz, mxxyPyzz, mxxyMyzz, mxxzPyyz, mxxzMyyz, mxyyPxzz, mxyyMxzz; //real qudricLimit = c1o100;//ganz schlechte Idee -> muss global sein - real O3 = c2o1 - o; + //real O3 = c2o1 - o; //real residu, residutmp; //residutmp = c0o1;///*-*/ c2o9 * (1./o - c1o2) * eps_new * eps_new; real NeqOn = c1o1;//zero;//one; //.... one = on ..... zero = off @@ -14484,7 +14484,7 @@ extern "C" __global__ void scaleFC_Fix_comp_27( real* DC, real m0, m1, m2, oMdrho; real mxxPyyPzz, mxxMyy, mxxMzz, mxxyPyzz, mxxyMyzz, mxxzPyyz, mxxzMyyz, mxyyPxzz, mxyyMxzz; //real qudricLimit = c1o100;//ganz schlechte Idee -> muss global sein - real O3 = c2o1 - o; + //real O3 = c2o1 - o; //real residu, residutmp; //residutmp = c0o1;///*-*/ c2o9 * (1./o - c1o2) * eps_new * eps_new; real NeqOn = c1o1;//zero;//one; //.... one = on ..... zero = off diff --git a/src/gpu/VirtualFluids_GPU/GPU/ScaleFC_F3_27.cu b/src/gpu/VirtualFluids_GPU/GPU/ScaleFC_F3_27.cu index b2437ce3f..3cccd8cb1 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/ScaleFC_F3_27.cu +++ b/src/gpu/VirtualFluids_GPU/GPU/ScaleFC_F3_27.cu @@ -801,7 +801,7 @@ extern "C" __global__ void scaleFC_comp_D3Q27F3_2018(real* DC, real m0, m1, m2, oMdrho; real mxxPyyPzz, mxxMyy, mxxMzz, mxxyPyzz, mxxyMyzz, mxxzPyyz, mxxzMyyz, mxyyPxzz, mxyyMxzz; //real qudricLimit = c1o100;//ganz schlechte Idee -> muss global sein - real O3 = c2o1 - o; + //real O3 = c2o1 - o; //real residu, residutmp; //residutmp = c0;///*-*/ c2o9 * (1./o - c1o2) * eps_new * eps_new; real NeqOn = c1o1;//zero;//one; //.... one = on ..... zero = off @@ -2048,7 +2048,7 @@ extern "C" __global__ void scaleFC_comp_D3Q27F3( real* DC, real m0, m1, m2, oMdrho; real mxxPyyPzz, mxxMyy, mxxMzz, mxxyPyzz, mxxyMyzz, mxxzPyyz, mxxzMyyz, mxyyPxzz, mxyyMxzz; //real qudricLimit = c1o100;//ganz schlechte Idee -> muss global sein - real O3 = c2o1 - o; + //real O3 = c2o1 - o; //real residu, residutmp; //residutmp = c0;///*-*/ c2o9 * (1./o - c1o2) * eps_new * eps_new; real NeqOn = c1o1;//zero;//one; //.... one = on ..... zero = off diff --git a/src/gpu/VirtualFluids_GPU/Input/PositionReader.cpp b/src/gpu/VirtualFluids_GPU/Input/PositionReader.cpp index ec91e92e5..b2a942c9a 100644 --- a/src/gpu/VirtualFluids_GPU/Input/PositionReader.cpp +++ b/src/gpu/VirtualFluids_GPU/Input/PositionReader.cpp @@ -91,7 +91,7 @@ const int INVDIR[] = { INV_E, static const int optionDigits = 2; //--> 2 bits f�r secondary Option static const long long maxOptionVal = ( 1<<optionDigits ) - 1; //2^3-1 -> 7 -float q[27]; +//float q[27]; long long noslipBoundaryFlags; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu index 246ec7965..814e239b7 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/Advection/Compressible/CumulantAll4/CumulantAll4CompSP27_Device.cu @@ -681,8 +681,8 @@ extern "C" __global__ void LB_Kernel_Cumulant_D3Q27All4( real omega, real dzuz = dxux + omega * c3o2 * mxxMzz; ////////////////////////////////////////////////////////////////////////////// - real divTest = (mxxPyyPzz - (mfaaa + (-c3o1 * (c1o1 - OxxPyyPzz*c1o2)*(vx2*dxux + vy2*dyuy + vz2*dzuz) + (c6o1 - c3o1 * (omega + OxxPyyPzz) + omega*OxxPyyPzz) / - (c3o1 * omega)*(dxux*dxux / rho + dyuy*dyuy / rho + dzuz*dzuz / rho + vvx*dxxux + vvy*dyyuy + vvz*dzzuz)) / OxxPyyPzz)) / OxxPyyPzz; + //real divTest = (mxxPyyPzz - (mfaaa + (-c3o1 * (c1o1 - OxxPyyPzz*c1o2)*(vx2*dxux + vy2*dyuy + vz2*dzuz) + (c6o1 - c3o1 * (omega + OxxPyyPzz) + omega*OxxPyyPzz) / + //(c3o1 * omega)*(dxux*dxux / rho + dyuy*dyuy / rho + dzuz*dzuz / rho + vvx*dxxux + vvy*dyyuy + vvz*dzzuz)) / OxxPyyPzz)) / OxxPyyPzz; //dxxux *= limAdvect / (limAdvect + sqrt(abs(divTest))); //dyyuy *= limAdvect / (limAdvect + sqrt(abs(divTest))); diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27_Device.cu index b36290b33..34c71db68 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Compressible/Mod27/ADComp27/ADComp27_Device.cu @@ -227,7 +227,7 @@ extern "C" __global__ void LB_KERNEL_AD_COMP_27(real diffusivity, (((mfbac + mfbca) + (mfbaa + mfbcc)) + ((mfabc + mfcba) + (mfaba + mfcbc)) + ((mfacb + mfcab) + (mfaab + mfccb))) + ((mfabb + mfcbb) + (mfbab + mfbcb)) + (mfbba + mfbbc)) + mfbbb; - real rho = c1o1 + drho; + //real rho = c1o1 + drho; //////////////////////////////////////////////////////////////////////////////// real rho0fluid = (fTNE + fBSW) + (fTSW + fBNE) + (fTSE + fBNW) + (fTNW + fBSE) + (fNE + fSW) + (fNW + fSE) + (fTE + fBW) + (fBE + fTW) + (fTN + fBS) + (fBN + fTS) + (fE + fW) + (fN + fS) + (fT + fB) + fZERO; real rhofluid = rho0fluid + c1o1; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7_Device.cu index f019f7c75..43d02e0f1 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/AdvectionDiffusion/Incompressible/Mod7/ADIncomp7/ADIncomp7_Device.cu @@ -197,9 +197,9 @@ extern "C" __global__ void LB_Kernel_AD_Incomp_7(real diffusivity, // real vy = ten * ((fTNE-fBSW)+(fBNE-fTSW)+(fBNW-fTSE)+(fTNW-fBSE) +(fNE-fSW)+(fNW-fSE)+(fTN-fBS)+(fBN-fTS)+(fN-fS)); // real vz = ten * ((fTNE-fBSW)+(fTSW-fBNE)+(fTSE-fBNW)+(fTNW-fBSE) +(fTE-fBW)+(fTW-fBE)+(fTN-fBS)+(fTS-fBN)+(fT-fB)); //////////////////////////////////////////////////////////////////////////////// - real ux_sq = vx * vx; - real uy_sq = vy * vy; - real uz_sq = vz * vz; + //real ux_sq = vx * vx; + //real uy_sq = vy * vy; + //real uz_sq = vz * vz; //////////////////////////////////////////////////////////////////////////////// //BGK //real omegaD = -three + sqrt(three); !!!!!!!!!!!!!!Achtung!!!!!!!!!!!!!!!!!! anderes Vorzeichen als in den Randbedingungen diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/Advection/Compressible/CumulantOne/PMCumulantOneCompSP27_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/Advection/Compressible/CumulantOne/PMCumulantOneCompSP27_Device.cu index 945d87064..734b71e7a 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/Advection/Compressible/CumulantOne/PMCumulantOneCompSP27_Device.cu +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/PorousMediaKernels/Advection/Compressible/CumulantOne/PMCumulantOneCompSP27_Device.cu @@ -475,7 +475,7 @@ extern "C" __global__ void LB_Kernel_PM_Cum_One_Comp_SP_27(real omega, real mxxMzz = mfcaa - mfaac; ////////////////////////////////////////////////////////////////////////// - real magicBulk = (CUMacc + CUMcac + CUMcca)*(c1o1 / OxxPyyPzz - c1o2)*c3o2*8.; + //real magicBulk = (CUMacc + CUMcac + CUMcca)*(c1o1 / OxxPyyPzz - c1o2)*c3o2*8.; ////////////////////////////////////////////////////////////////////////// //limiter-Scheise Teil 1 diff --git a/src/gpu/VirtualFluids_GPU/Output/DataWriter.h b/src/gpu/VirtualFluids_GPU/Output/DataWriter.h index 5e7ee8b30..6e61fc696 100644 --- a/src/gpu/VirtualFluids_GPU/Output/DataWriter.h +++ b/src/gpu/VirtualFluids_GPU/Output/DataWriter.h @@ -22,7 +22,5 @@ public: virtual void writeInit(std::shared_ptr<Parameter> para, std::shared_ptr<CudaMemoryManager> cudaManager) = 0; virtual void writeTimestep(std::shared_ptr<Parameter> para, unsigned int timestep) = 0; virtual void writeTimestep(std::shared_ptr<Parameter> para, unsigned int timestep, int level) = 0; - - DataWriter(const DataWriter& dataWriter) {} }; #endif diff --git a/src/gpu/VirtualFluids_GPU/Output/LogWriter.hpp b/src/gpu/VirtualFluids_GPU/Output/LogWriter.hpp index 587364be9..cbce9a48c 100644 --- a/src/gpu/VirtualFluids_GPU/Output/LogWriter.hpp +++ b/src/gpu/VirtualFluids_GPU/Output/LogWriter.hpp @@ -21,9 +21,9 @@ public: consoleOut = false; this->fname = fname; } - void setName(std::string fname) + void setName(std::string name) { - this->fname = fname; + this->fname = name; } void setConsoleOut(bool flag) { diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp index d23d54265..2145644c6 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp @@ -395,9 +395,9 @@ Parameter::Parameter(SPtr<ConfigData> configData, Communicator* comm) this->setForcing(forcingX, forcingY, forcingZ); ////////////////////////////////////////////////////////////////////////// //quadricLimiters - real quadricLimiterP = 0.01; - real quadricLimiterM = 0.01; - real quadricLimiterD = 0.01; + real quadricLimiterP = (real)0.01; + real quadricLimiterM = (real)0.01; + real quadricLimiterD = (real)0.01; if (configData->isQuadricLimiterPInConfigFile()) quadricLimiterP = configData->getQuadricLimiterP(); -- GitLab