diff --git a/gpu.cmake b/gpu.cmake index e31d81de64dde6c206372e0f5cda99cf450f848f..6c0f3d489c668b029417eb9147966db13acd643d 100644 --- a/gpu.cmake +++ b/gpu.cmake @@ -43,7 +43,7 @@ IF (BUILD_VF_GPU) #add_subdirectory(apps/gpu/LBM/gridGeneratorTest) #add_subdirectory(apps/gpu/LBM/TGV_3D) #add_subdirectory(apps/gpu/LBM/TGV_3D_MultiGPU) - #add_subdirectory(apps/gpu/LBM/ActuatorLine) + add_subdirectory(apps/gpu/LBM/ActuatorLine) ELSE() MESSAGE( STATUS "exclude Virtual Fluids GPU." ) ENDIF() diff --git a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu index 6a2a625f03e64059f59480a3938f606d9f63d254..a3994dc00cee1f8cf028f5ddaaacce515e6414f9 100644 --- a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu +++ b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu @@ -26,10 +26,10 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real uint* neighborsX, uint* neighborsY, uint* neighborsZ, uint* neighborsWSB, real* vx, real* vy, real* vz, - int numberOfIndices, + uint numberOfIndices, real* bladeCoordsX, real* bladeCoordsY, real* bladeCoordsZ, real* bladeVelocitiesX, real* bladeVelocitiesY, real* bladeVelocitiesZ, - uint* bladeIndices, int numberOfNodes) + uint* bladeIndices, uint numberOfNodes) { // Possibly restructure this to check every cell whether it is bsw of some blade node and then do interpolation // --> no need to save closest nodes and search new closest node @@ -89,12 +89,12 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real __global__ void applyBodyForces(real* gridCoordsX, real* gridCoordsY, real* gridCoordsZ, real* gridForcesX, real* gridForcesY, real* gridForcesZ, - uint* gridIndices, int numberOfIndices, + uint* gridIndices, uint numberOfIndices, real* bladeCoordsX, real* bladeCoordsY, real* bladeCoordsZ, real* bladeForcesX, real* bladeForcesY,real* bladeForcesZ, real* bladeRadii, real radius, - int nBlades, int nBladeNodes, + uint nBlades, uint nBladeNodes, real epsilon, real delta_x) { const uint x = threadIdx.x; @@ -253,7 +253,7 @@ void ActuatorLine::calcForcesEllipticWing(Parameter* para) real forceRatio = this->density*pow(this->delta_x,4)*pow(this->delta_t,-2); - for( int blade=0; blade<this->nBlades; blade++) + for( uint blade=0; blade<this->nBlades; blade++) { localAzimuth = this->azimuth+2*blade*vf::lbm::constant::cPi/this->nBlades; for( uint bladeNode=0; bladeNode<this->nBladeNodes; bladeNode++) @@ -312,7 +312,7 @@ void ActuatorLine::calcForcesEllipticWing(Parameter* para) void ActuatorLine::rotateBlades(real angle) { - for(unsigned int node=0; node<this->nBladeNodes*this->nBlades; node++) + for(uint node=0; node<this->nBladeNodes*this->nBlades; node++) { real oldCoordX = this->bladeCoordsXH[node]; real oldCoordY = this->bladeCoordsYH[node]; @@ -343,10 +343,10 @@ void ActuatorLine::initBladeCoords(CudaMemoryManager* cudaManager) { cudaManager->cudaAllocBladeCoords(this); - for( unsigned int blade=0; blade<this->nBlades; blade++) + for( uint blade=0; blade<this->nBlades; blade++) { real localAzimuth = this->azimuth+(2*vf::lbm::constant::cPi/this->nBlades)*blade; - for(unsigned int node=0; node<this->nBladeNodes; node++) + for(uint node=0; node<this->nBladeNodes; node++) { real coordX, coordY, coordZ; real x,y,z; @@ -367,7 +367,7 @@ void ActuatorLine::initBladeVelocities(CudaMemoryManager* cudaManager) { cudaManager->cudaAllocBladeVelocities(this); - for(unsigned int node=0; node<this->numberOfNodes; node++) + for(uint node=0; node<this->numberOfNodes; node++) { this->bladeVelocitiesXH[node] = 0.f; this->bladeVelocitiesYH[node] = 0.f; @@ -380,7 +380,7 @@ void ActuatorLine::initBladeForces(CudaMemoryManager* cudaManager) { cudaManager->cudaAllocBladeForces(this); - for(unsigned int node=0; node<this->numberOfNodes; node++) + for(uint node=0; node<this->numberOfNodes; node++) { this->bladeForcesXH[node] = 0.f; this->bladeForcesYH[node] = 0.f; @@ -397,7 +397,7 @@ void ActuatorLine::initBladeIndices(Parameter* para, CudaMemoryManager* cudaMana real* coordsY = para->getParH(this->level)->coordY_SP; real* coordsZ = para->getParH(this->level)->coordZ_SP; - for(unsigned int node=0; node<this->numberOfNodes; node++) + for(uint node=0; node<this->numberOfNodes; node++) { this->bladeIndicesH[node] = findNearestCellBSW(1, coordsX, coordsY, coordsZ, this->bladeCoordsXH[node], this->bladeCoordsYH[node], this->bladeCoordsZH[node], @@ -413,7 +413,7 @@ void ActuatorLine::initBoundingSphere(Parameter* para, CudaMemoryManager* cudaMa // Actuator line exists only on 1 level std::vector<int> nodesInSphere; - for (int j = 1; j <= para->getParH(this->level)->size_Mat_SP; j++) + for (uint j = 1; j <= para->getParH(this->level)->size_Mat_SP; j++) { const real coordX = para->getParH(this->level)->coordX_SP[j]; const real coordY = para->getParH(this->level)->coordY_SP[j]; diff --git a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h index 4c5a55d01cc632b2456c16d315f49886506568ff..6863840e3082b980a081a98d24a7a1b49fb79ba7 100644 --- a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h +++ b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h @@ -10,9 +10,9 @@ class ActuatorLine : public Visitor { public: ActuatorLine( - const unsigned int _nBlades, + const uint _nBlades, const real _density, - const unsigned int _nBladeNodes, + const uint _nBladeNodes, const real _epsilon, real _turbinePosX, real _turbinePosY, real _turbinePosZ, const real _diameter, @@ -43,7 +43,7 @@ public: } void init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaManager); - void visit(Parameter* para, CudaMemoryManager* cudaManager, int level, unsigned int t); + void visit(Parameter* para, CudaMemoryManager* cudaManager, int level, uint t); void free(Parameter* para, CudaMemoryManager* cudaManager); void write(uint t); @@ -88,12 +88,12 @@ private: real turbinePosX, turbinePosY, turbinePosZ; real omega, azimuth, delta_t, delta_x; const real diameter; - const unsigned int nBladeNodes; - const unsigned int nBlades; + const uint nBladeNodes; + const uint nBlades; const real epsilon; // in m const int level; - int numberOfIndices; - int numberOfNodes; + uint numberOfIndices; + uint numberOfNodes; }; #endif \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Visitor/Probe.cu b/src/gpu/VirtualFluids_GPU/Visitor/Probe.cu index 31e082c8640225fcb2b790f2f6dec19c78c5f3ec..5d317082b27206b90ec48ed7612a3effbf565567 100644 --- a/src/gpu/VirtualFluids_GPU/Visitor/Probe.cu +++ b/src/gpu/VirtualFluids_GPU/Visitor/Probe.cu @@ -15,14 +15,14 @@ #include "GPU/CudaMemoryManager.h" -__global__ void interpQuantities( int* pointIndices, +__global__ void interpQuantities( uint* pointIndices, uint nPoints, real* distX, real* distY, real* distZ, real* vx, real* vy, real* vz, real* rho, uint* neighborX, uint* neighborY, uint* neighborZ, // real* vx_point, real* vy_point, real* vz_point, real* rho_point, PostProcessingVariable* PostProcessingVariables, - int* quantityArrayOffsets, real* quantityArray + uint* quantityArrayOffsets, real* quantityArray ) { const uint x = threadIdx.x; @@ -64,7 +64,7 @@ __global__ void interpQuantities( int* pointIndices, for( int variableIndex = 0; PostProcessingVariables[variableIndex] != PostProcessingVariable::LAST; variableIndex++) { PostProcessingVariable variable = PostProcessingVariables[variableIndex]; - int arrayOffset = quantityArrayOffsets[int(variable)]*nPoints; + uint arrayOffset = quantityArrayOffsets[int(variable)]*nPoints; switch(variable) { @@ -153,7 +153,7 @@ void Probe::init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaManager->cudaCopyProbeDistancesHtoD(this, level); cudaManager->cudaCopyProbeIndicesHtoD(this, level); - int arrOffset = 0; + uint arrOffset = 0; cudaManager->cudaAllocProbeQuantities(this, level); @@ -265,11 +265,11 @@ void Probe::addPostProcessingVariable(PostProcessingVariable _variable) void Probe::write(Parameter* para, int level, int t) { - const unsigned int numberOfParts = this->getProbeStruct(level)->nPoints / para->getlimitOfNodesForVTK() + 1; + const uint numberOfParts = this->getProbeStruct(level)->nPoints / para->getlimitOfNodesForVTK() + 1; std::vector<std::string> fnames; - for (unsigned int i = 1; i <= numberOfParts; i++) + for (uint i = 1; i <= numberOfParts; i++) { fnames.push_back(this->probeName + "_bin_lev_" + StringUtil::toString<int>(level) + "_ID_" + StringUtil::toString<int>(para->getMyID()) + "_Part_" + StringUtil::toString<int>(i) + "_t_" + StringUtil::toString<int>(t) + ".vtk"); this->fileNamesForCollectionFile.push_back(fnames.back()); @@ -328,13 +328,13 @@ void Probe::writeGridFile(Parameter* para, int level, std::vector<std::string>& std::vector< UbTupleFloat3 > nodes; std::vector< std::string > nodedatanames = this->getVarNames(); - unsigned int startpos = 0; - unsigned int endpos = 0; - unsigned int sizeOfNodes = 0; + uint startpos = 0; + uint endpos = 0; + uint sizeOfNodes = 0; std::vector< std::vector< double > > nodedata(nodedatanames.size()); real inv_t = 1/(real(max(t,1))*pow(2,level)); - for (unsigned int part = 0; part < fnames.size(); part++) + for (uint part = 0; part < fnames.size(); part++) { startpos = part * para->getlimitOfNodesForVTK(); sizeOfNodes = min(para->getlimitOfNodesForVTK(), this->getProbeStruct(level)->nPoints - startpos); @@ -410,4 +410,3 @@ std::vector<std::string> Probe::getVarNames() } return varNames; } - \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Visitor/Probe.h b/src/gpu/VirtualFluids_GPU/Visitor/Probe.h index 98a00c640da11e97c325b891d6615cf574f8ed98..5f5c7ea8266f1305b988211523dc826def088720 100644 --- a/src/gpu/VirtualFluids_GPU/Visitor/Probe.h +++ b/src/gpu/VirtualFluids_GPU/Visitor/Probe.h @@ -17,13 +17,13 @@ enum class PostProcessingVariable{ LAST = 8, }; struct ProbeStruct{ - int nPoints, nArrays; - int *pointIndicesH, *pointIndicesD; + uint nPoints, nArrays; + uint *pointIndicesH, *pointIndicesD; real *pointCoordsX, *pointCoordsY, *pointCoordsZ; real *distXH, *distYH, *distZH, *distXD, *distYD, *distZD; real* quantitiesArrayH, *quantitiesArrayD; PostProcessingVariable* quantitiesH,* quantitiesD; - int* arrayOffsetsH, *arrayOffsetsD; + uint* arrayOffsetsH, *arrayOffsetsD; };