From 4bf413805f3d9bed1cadb036a2955202e3091a07 Mon Sep 17 00:00:00 2001 From: Hkorb <henry.korb@geo.uu.se> Date: Thu, 30 Sep 2021 18:25:42 +0200 Subject: [PATCH] minor changes to ActuatorLine and Probes changes raw pointer to shared pointer --- .../Calculation/UpdateGrid27.cpp | 6 +-- src/gpu/VirtualFluids_GPU/GPU/GeometryUtils.h | 49 ++++++++++--------- src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 8 +-- .../VirtualFluids_GPU/Parameter/Parameter.cpp | 8 +-- .../VirtualFluids_GPU/Parameter/Parameter.h | 12 ++--- .../PreCollisionInteractor/ActuatorLine.cu | 5 +- .../PreCollisionInteractor/ActuatorLine.h | 5 +- .../Probes/PlaneProbe.cu | 2 +- .../Probes/PointProbe.cu | 2 +- .../PreCollisionInteractor/Probes/Probe.cu | 2 +- 10 files changed, 50 insertions(+), 49 deletions(-) diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp index 8fed82906..27759453e 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 63473780c..da438ab0f 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 a7e42e41e..78f8742b9 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 86b8f1d4f..2aa00708e 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 4b30d9ff9..06161934b 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 ab2de3cab..1f800b1bd 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 cb8553988..b2fad77f1 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 440f60d81..e73c732f7 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 7de22cc18..40faab4ee 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 063792712..b1a31f2b2 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; -- GitLab