diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp index 8fed82906fcb0d6b6547cd17286fa5ea9bc6f601..27759453e424bc35ef0f73a354c82415dfb53331 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp +++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp @@ -63,7 +63,7 @@ void updateGrid27(Parameter* para, coarseToFine(para, level); } - visitPreCollisionInteractors(para, cudaManager, level, t); + visitActuators(para, cudaManager, level, t); visitProbes(para, cudaManager, level, t); } @@ -1269,7 +1269,7 @@ void coarseToFine(Parameter* para, int level) void visitActuators(Parameter* para, CudaMemoryManager* cudaManager, int level, unsigned int t) { - for( PreCollisionInteractor* actuator: para->getActuators() ) + for( SPtr<PreCollisionInteractor> actuator: para->getActuators() ) { actuator->visit(para, cudaManager, level, t); } @@ -1277,7 +1277,7 @@ void visitActuators(Parameter* para, CudaMemoryManager* cudaManager, int level, void visitProbes(Parameter* para, CudaMemoryManager* cudaManager, int level, unsigned int t) { - for( PreCollisionInteractor* probe: para->getProbes() ) + for( SPtr<PreCollisionInteractor> probe: para->getProbes() ) { probe->visit(para, cudaManager, level, t); } diff --git a/src/gpu/VirtualFluids_GPU/GPU/GeometryUtils.h b/src/gpu/VirtualFluids_GPU/GPU/GeometryUtils.h index 63473780cf0548b0a8ffa97a0d402497030de3fb..da438ab0fdb04c8a83bd37700b4e4735970bcd7d 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/GeometryUtils.h +++ b/src/gpu/VirtualFluids_GPU/GPU/GeometryUtils.h @@ -1,7 +1,7 @@ #ifndef _GEOMETRYUTILS_H #define _GEOMETRYUTILS_H -__inline__ __host__ __device__ void getNeighborIndicesBSW( uint k, //index of BSW node +__inline__ __host__ __device__ void getNeighborIndicesOfBSW( uint k, //index of BSW node uint &ke, uint &kn, uint &kt, uint &kne, uint &kte,uint &ktn, uint &ktne, uint* neighborX, uint* neighborY, uint* neighborZ) { @@ -14,6 +14,30 @@ __inline__ __host__ __device__ void getNeighborIndicesBSW( uint k, //index of B ktne = neighborX[ktn]; } +__inline__ __host__ __device__ uint findNearestCellBSW(uint index, + real* coordsX, real* coordsY, real* coordsZ, + real posX, real posY, real posZ, + uint* neighborsX, uint* neighborsY, uint* neighborsZ, uint* neighborsWSB) +{ + uint new_index = index; + + while(coordsX[new_index] > posX && coordsY[new_index] > posY && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsWSB[new_index]);} + + while(coordsX[new_index] > posX && coordsY[new_index] > posY ){ new_index = max(1, neighborsZ[neighborsWSB[new_index]]);} + while(coordsX[new_index] > posX && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsY[neighborsWSB[new_index]]);} + while(coordsY[new_index] > posY && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsX[neighborsWSB[new_index]]);} + + while(coordsX[new_index] > posX){ new_index = max(1, neighborsY[neighborsZ[neighborsWSB[new_index]]]);} + while(coordsY[new_index] > posY){ new_index = max(1, neighborsX[neighborsZ[neighborsWSB[new_index]]]);} + while(coordsZ[new_index] > posZ){ new_index = max(1, neighborsX[neighborsY[neighborsWSB[new_index]]]);} + + while(coordsX[new_index] < posX){ new_index = max(1, neighborsX[new_index]);} + while(coordsY[new_index] < posY){ new_index = max(1, neighborsY[new_index]);} + while(coordsZ[new_index] < posZ){ new_index = max(1, neighborsZ[new_index]);} + + return neighborsWSB[new_index]; +} + __inline__ __host__ __device__ void getInterpolationWeights(real &dW, real &dE, real &dN, real &dS, real &dT, real &dB, real tmpX, real tmpY, real tmpZ) { @@ -175,28 +199,5 @@ __inline__ __host__ __device__ void invRotateAboutZ3D(real &angle, real &posX, r translate3D(tmpX, tmpY, tmpZ, newPosX, newPosY, newPosZ, originX, originY, originZ); } -__inline__ __host__ __device__ uint findNearestCellBSW(uint index, - real* coordsX, real* coordsY, real* coordsZ, - real posX, real posY, real posZ, - uint* neighborsX, uint* neighborsY, uint* neighborsZ, uint* neighborsWSB) -{ - uint new_index = index; - - while(coordsX[new_index] > posX && coordsY[new_index] > posY && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsWSB[new_index]);} - - while(coordsX[new_index] > posX && coordsY[new_index] > posY ){ new_index = max(1, neighborsZ[neighborsWSB[new_index]]);} - while(coordsX[new_index] > posX && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsY[neighborsWSB[new_index]]);} - while(coordsY[new_index] > posY && coordsZ[new_index] > posZ ){ new_index = max(1, neighborsX[neighborsWSB[new_index]]);} - - while(coordsX[new_index] > posX){ new_index = max(1, neighborsY[neighborsZ[neighborsWSB[new_index]]]);} - while(coordsY[new_index] > posY){ new_index = max(1, neighborsX[neighborsZ[neighborsWSB[new_index]]]);} - while(coordsZ[new_index] > posZ){ new_index = max(1, neighborsX[neighborsY[neighborsWSB[new_index]]]);} - - while(coordsX[new_index] < posX){ new_index = max(1, neighborsX[new_index]);} - while(coordsY[new_index] < posY){ new_index = max(1, neighborsY[new_index]);} - while(coordsZ[new_index] < posZ){ new_index = max(1, neighborsZ[new_index]);} - - return neighborsWSB[new_index]; -} #endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp index a7e42e41ef4f65f419b52bec6cbb108d5270c6b7..78f8742b915f1465a35da47c2c560cd32693ce36 100644 --- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp +++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp @@ -130,11 +130,11 @@ void Simulation::init(SPtr<Parameter> para, SPtr<GridProvider> gridProvider, std gridProvider->allocArrays_BoundaryQs(); gridProvider->allocArrays_OffsetScale(); - for( PreCollisionInteractor* actuator: para->getActuators()){ + for( SPtr<PreCollisionInteractor> actuator: para->getActuators()){ actuator->init(para.get(), gridProvider.get(), cudaManager.get()); } - for( PreCollisionInteractor* probe: para->getProbes()){ + for( SPtr<PreCollisionInteractor> probe: para->getProbes()){ probe->init(para.get(), gridProvider.get(), cudaManager.get()); } @@ -1288,11 +1288,11 @@ void Simulation::free() } ////////////////////////////////////////////////////////////////////////// //PreCollisionInteractors - for( PreCollisionInteractor* actuator: para->getActuators()){ + for( SPtr<PreCollisionInteractor> actuator: para->getActuators()){ actuator->free(para.get(), cudaManager.get()); } - for( PreCollisionInteractor* probe: para->getProbes()){ + for( SPtr<PreCollisionInteractor> probe: para->getProbes()){ probe->free(para.get(), cudaManager.get()); } ////////////////////////////////////////////////////////////////////////// diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp index 86b8f1d4fccbeb999f081a8cbc0ea9c47b13ef9a..2aa00708e5a440f99fe341374dd668a5e764bbdb 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp @@ -1492,11 +1492,11 @@ void Parameter::setADKernel(std::string adKernel) //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// //add-methods //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// -void Parameter::addActuator(PreCollisionInteractor* actuator) +void Parameter::addActuator(SPtr<PreCollisionInteractor> actuator) { actuators.push_back(actuator); } -void Parameter::addProbe(PreCollisionInteractor* probe) +void Parameter::addProbe(SPtr<PreCollisionInteractor> probe) { probes.push_back(probe); } @@ -1911,11 +1911,11 @@ TempPressforBoundaryConditions* Parameter::getTempPressD() { return this->TempPressD; } -std::vector<PreCollisionInteractor*> Parameter::getActuators() +std::vector<SPtr<PreCollisionInteractor>> Parameter::getActuators() { return actuators; } -std::vector<PreCollisionInteractor*> Parameter::getProbes() +std::vector<SPtr<PreCollisionInteractor>> Parameter::getProbes() { return probes; } diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h index 4b30d9ff9b8bd6b945cade5b79870458c3947997..06161934b52e35c2a0932277a651bbdc102424a8 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h @@ -526,8 +526,8 @@ public: //adder - void addActuator(PreCollisionInteractor* actuator); - void addProbe(PreCollisionInteractor* probes); + void addActuator(SPtr<PreCollisionInteractor> actuator); + void addProbe(SPtr<PreCollisionInteractor> probes); // getter double *getForcesDouble(); @@ -687,8 +687,8 @@ public: TempVelforBoundaryConditions *getTempVelD(); TempPressforBoundaryConditions *getTempPressH(); TempPressforBoundaryConditions *getTempPressD(); - std::vector<PreCollisionInteractor*> getActuators(); - std::vector<PreCollisionInteractor*> getProbes(); + std::vector<SPtr<PreCollisionInteractor>> getActuators(); + std::vector<SPtr<PreCollisionInteractor>> getProbes(); unsigned int getTimeDoCheckPoint(); unsigned int getTimeDoRestart(); bool getDoCheckPoint(); @@ -820,8 +820,8 @@ private: unsigned int startTurn; // PreCollisionInteractors ////////////// - std::vector<PreCollisionInteractor*> actuators; - std::vector<PreCollisionInteractor*> probes; + std::vector<SPtr<PreCollisionInteractor>> actuators; + std::vector<SPtr<PreCollisionInteractor>> probes; // Step of Ensight writing// unsigned int stepEnsight; diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu index ab2de3cabe83448019f5a78b55c9fdf931df2fb9..1f800b1bd8cd3e5520693b2460988f1f16fa6778 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.cu @@ -58,7 +58,7 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real bladeIndices[node] = k; - getNeighborIndicesBSW(k, ke, kn, kt, kne, kte, ktn, ktne, neighborsX, neighborsY, neighborsZ); + getNeighborIndicesOfBSW(k, ke, kn, kt, kne, kte, ktn, ktne, neighborsX, neighborsY, neighborsZ); real dW, dE, dN, dS, dT, dB; @@ -158,7 +158,7 @@ void ActuatorLine::visit(Parameter* para, CudaMemoryManager* cudaManager, int le cudaManager->cudaCopyBladeCoordsHtoD(this); - uint numberOfThreads = 128; + uint numberOfThreads = para->getParH(level)->numberofthreads; vf::gpu::CudaGrid bladeGrid = vf::gpu::CudaGrid(numberOfThreads, this->numberOfNodes); interpolateVelocities<<< bladeGrid.grid, bladeGrid.threads >>>( @@ -206,7 +206,6 @@ void ActuatorLine::free(Parameter* para, CudaMemoryManager* cudaManager) cudaManager->cudaFreeBladeVelocities(this); cudaManager->cudaFreeBladeForces(this); cudaManager->cudaFreeBladeIndices(this); - cudaManager->cudaFreeSphereIndices(this); } diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h index cb85539886c43701a3587ae055b6485f4a972d5f..b2fad77f1d4a1e494c463a9e1571b5b641faa638 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorLine.h @@ -2,9 +2,10 @@ #define ActuatorLine_H #include "PreCollisionInteractor.h" -#include "Parameter/Parameter.h" #include "PointerDefinitions.h" -#include "GridGenerator/grid/GridBuilder/GridBuilder.h" + +class Parameter; +class GridProvider; class ActuatorLine : public PreCollisionInteractor { diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu index 440f60d81bd1618aad9e4e64ebfe76fc6568a22d..e73c732f7fa3a592c04d4406472ffabdee71ed4d 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PlaneProbe.cu @@ -41,7 +41,7 @@ void PlaneProbe::findPoints(Parameter* para, GridProvider* gridProvider, std::ve void PlaneProbe::calculateQuantities(ProbeStruct* probeStruct, Parameter* para, int level) { - vf::gpu::CudaGrid grid = vf::gpu::CudaGrid(128, probeStruct->nPoints); + vf::gpu::CudaGrid grid = vf::gpu::CudaGrid(para->getParH(level)->numberofthreads, probeStruct->nPoints); interpQuantities<<<grid.grid, grid.threads>>>( probeStruct->pointIndicesD, probeStruct->nPoints, probeStruct->vals, probeStruct->distXD, probeStruct->distYD, probeStruct->distZD, para->getParD(level)->vx_SP, para->getParD(level)->vy_SP, para->getParD(level)->vz_SP, para->getParD(level)->rho_SP, diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu index 7de22cc181f5b2a75b9d522b82c3c4336e562d91..40faab4eeaab38c04ec65bf6a8ae2679fed24fe6 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/PointProbe.cu @@ -44,7 +44,7 @@ void PointProbe::findPoints(Parameter* para, GridProvider* gridProvider, std::ve void PointProbe::calculateQuantities(ProbeStruct* probeStruct, Parameter* para, int level) { - vf::gpu::CudaGrid grid = vf::gpu::CudaGrid(128, probeStruct->nPoints); + vf::gpu::CudaGrid grid = vf::gpu::CudaGrid(para->getParH(level)->numberofthreads, probeStruct->nPoints); interpQuantities<<<grid.grid, grid.threads>>>( probeStruct->pointIndicesD, probeStruct->nPoints, probeStruct->vals, probeStruct->distXD, probeStruct->distYD, probeStruct->distZD, diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu index 063792712b2b0f3173499179f77e5340ea63e090..b1a31f2b2df0e2c016c1c327807420f1b0f267ea 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/Probes/Probe.cu @@ -112,7 +112,7 @@ __global__ void interpQuantities( uint* pointIndices, if(interpolate) { uint ke, kn, kt, kne, kte, ktn, ktne; - getNeighborIndicesBSW( k, ke, kn, kt, kne, kte, ktn, ktne, neighborX, neighborY, neighborZ); + getNeighborIndicesOfBSW( k, ke, kn, kt, kne, kte, ktn, ktne, neighborX, neighborY, neighborZ); // Trilinear interpolation of macroscopic quantities to probe point real dW, dE, dN, dS, dT, dB;