From d0ff1af9f970010dee109995dce594a6a92848a3 Mon Sep 17 00:00:00 2001
From: Henry Korb <henry.korb@geo.uu.se>
Date: Tue, 22 Jun 2021 18:20:11 +0200
Subject: [PATCH] added visitors, added actuator_line
---
apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp | 16 +-
.../Calculation/UpdateGrid27.cpp | 20 ++
.../Calculation/UpdateGrid27.h | 5 +
.../DataStructureInitializer/GridProvider.cpp | 6 +
.../GridReaderGenerator/GridGenerator.cpp | 5 +
.../GPU/CudaMemoryManager.cpp | 2 +-
.../CumulantK17CompChim_Device.cu | 5 +
src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 7 +
.../VirtualFluids_GPU/Parameter/Parameter.cpp | 23 +++
.../VirtualFluids_GPU/Parameter/Parameter.h | 15 ++
.../VirtualFluids_GPU/Visitor/ActuatorLine.cu | 177 ++++++++++++++++++
.../VirtualFluids_GPU/Visitor/ActuatorLine.h | 78 ++++++++
src/gpu/VirtualFluids_GPU/Visitor/Visitor.h | 58 ++++++
13 files changed, 413 insertions(+), 4 deletions(-)
create mode 100644 src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu
create mode 100644 src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h
create mode 100644 src/gpu/VirtualFluids_GPU/Visitor/Visitor.h
diff --git a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
index d1d49e089..e225998ab 100644
--- a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
+++ b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp
@@ -45,6 +45,7 @@
#include "VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.h"
#include "VirtualFluids_GPU/Parameter/Parameter.h"
#include "VirtualFluids_GPU/Output/FileWriter.h"
+#include "VirtualFluids_GPU/Visitor/ActuatorLine.h"
#include "VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.h"
#include "VirtualFluids_GPU/PreProcessor/PreProcessorFactory/PreProcessorFactoryImp.h"
@@ -98,7 +99,7 @@ const real velocity = 9.0;
const real mach = 0.1;
-const uint nodes_per_D = 32;
+const uint nodes_per_D = 8;
//std::string path("F:/Work/Computations/out/DrivenCavity/"); //LEGOLAS
//std::string path("D:/out/DrivenCavity"); //Mollok
@@ -202,7 +203,8 @@ void multipleLevel(const std::string& configPath)
para->setTOut( timeStepOut );
para->setTEnd( timeStepEnd );
- para->setIsBodyForce( false );
+ para->setIsBodyForce( true );
+
/////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
gridBuilder->setVelocityBoundaryCondition(SideType::MX, velocityLB, 0.0, 0.0);
@@ -218,12 +220,20 @@ void multipleLevel(const std::string& configPath)
SPtr<GridProvider> gridGenerator = GridProvider::makeGridGenerator(gridBuilder, para, cudaMemoryManager);
+ real turbPos[3] = {100.f, 100.f, 100.f};
+ real epsilon = 5.f;
+ real density = 1.225f;
+ int level = 0;
+
+ ActuatorLine* actuator_line = new ActuatorLine((unsigned int) 3, density, (unsigned int)32, epsilon, turbPos[0], turbPos[1], turbPos[2], D, level);
+ para->addActuator( actuator_line );
+
Simulation sim;
SPtr<FileWriter> fileWriter = SPtr<FileWriter>(new FileWriter());
SPtr<KernelFactoryImp> kernelFactory = KernelFactoryImp::getInstance();
SPtr<PreProcessorFactoryImp> preProcessorFactory = PreProcessorFactoryImp::getInstance();
sim.setFactories(kernelFactory, preProcessorFactory);
- sim.init(para, gridGenerator, fileWriter, cudaMemoryManager);
+ sim.init(para, gridGenerator, fileWriter, cudaMemoryManager);
sim.run();
sim.free();
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index 9b90e0364..229bcc444 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -59,6 +59,10 @@ void updateGrid27(Parameter* para,
coarseToFine(para, level);
}
+
+ visitVisitors(para, level, t);
+
+ visitProbes(para, level, t);
}
void collision(Parameter* para, std::vector<std::shared_ptr<PorousMedia>>& pm, int level, unsigned int t, std::vector < SPtr< Kernel>>& kernels)
@@ -1259,3 +1263,19 @@ void coarseToFine(Parameter* para, int level)
}
}
+
+void visitVisitors(Parameter* para, int level, unsigned int t)
+{
+ for( Visitor* actuator: para->getActuators() )
+ {
+ actuator->visit(para, level, t);
+ }
+}
+
+void visitProbes(Parameter* para, int level, unsigned int t)
+{
+ for( Visitor* probe: para->getProbes() )
+ {
+ probe->visit(para, level, t);
+ }
+}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index c66d6afd4..1a3f39484 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -38,4 +38,9 @@ extern "C" void fineToCoarse(Parameter* para, int level);
extern "C" void coarseToFine(Parameter* para, int level);
+extern "C" void visitVisitors(Parameter* para, int level, unsigned int t);
+
+extern "C" void visitProbes(Parameter* para, int level, unsigned int t);
+
+
#endif
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index 957935756..f21ee67a1 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -81,6 +81,12 @@ void GridProvider::setInitalNodeValues(const int numberOfNodes, const int level)
para->getParH(level)->gDyvz[j] = 0.0f;
para->getParH(level)->gDzvz[j] = 0.0f;
}
+
+ if (para->getIsBodyForce()) {
+ para->getParH(level)->forceX_SP[j] = 0.0f;
+ para->getParH(level)->forceY_SP[j] = 0.0f;
+ para->getParH(level)->forceZ_SP[j] = 0.0f;
+ }
}
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index 3ce1093d1..6edd22963 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -63,6 +63,9 @@ void GridGenerator::allocArrays_CoordNeighborGeo()
if(para->getUseWale())
cudaMemoryManager->cudaAllocTurbulentViscosity(level);
+
+ if(para->getIsBodyForce())
+ cudaMemoryManager->cudaAllocBodyForce(level);
builder->getNodeValues(
para->getParH(level)->coordX_SP,
@@ -80,6 +83,8 @@ void GridGenerator::allocArrays_CoordNeighborGeo()
cudaMemoryManager->cudaCopyNeighborWSB(level);
cudaMemoryManager->cudaCopySP(level);
cudaMemoryManager->cudaCopyCoord(level);
+ if(para->getIsBodyForce())
+ cudaMemoryManager->cudaCopyBodyForce(level);
//std::cout << verifyNeighborIndices(level);
}
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 035ee8f80..a8df07396 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -40,7 +40,7 @@ void CudaMemoryManager::cudaAllocCoord(int lev)
checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordX_SP ), parameter->getParH(lev)->mem_size_real_SP ));
checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordY_SP ), parameter->getParH(lev)->mem_size_real_SP ));
checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordZ_SP ), parameter->getParH(lev)->mem_size_real_SP ));
- //Device (spinning ship + upsala)
+ //Device (spinning ship + uppsala)
checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordX_SP ), parameter->getParH(lev)->mem_size_real_SP ));
checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordY_SP ), parameter->getParH(lev)->mem_size_real_SP ));
checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordZ_SP ), parameter->getParH(lev)->mem_size_real_SP ));
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim_Device.cu
index 6676099ad..2b26cd2b6 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim_Device.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim_Device.cu
@@ -223,6 +223,11 @@ extern "C" __global__ void LB_Kernel_CumulantK17CompChim(
fx += bodyForceX[k];
fy += bodyForceY[k];
fz += bodyForceZ[k];
+
+ //Reset body force
+ bodyForceX[k] = 0.0f;
+ bodyForceY[k] = 0.0f;
+ bodyForceZ[k] = 0.0f;
}
vvx += fx * c1o2 / factor;
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 7dc916a47..47b57daca 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -140,6 +140,13 @@ void Simulation::init(SPtr<Parameter> para, SPtr<GridProvider> gridProvider, std
gridProvider->allocArrays_BoundaryQs();
gridProvider->allocArrays_OffsetScale();
+ for( Visitor* actuator: para->getActuators()){
+ actuator->init(para.get(), gridProvider.get(), cudaManager.get());
+ }
+
+ for( Visitor* probe: para->getProbes()){
+ probe->init(para.get(), gridProvider.get(), cudaManager.get());
+ }
//////////////////////////////////////////////////////////////////////////
//Kernel init
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
index 0446608cf..e4549b2cd 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
@@ -18,6 +18,8 @@
#include "Core/Input/ConfigData/ConfigData.h"
#include "Core/StringUtilities/StringUtil.h"
#include "Communication/Communicator.h"
+#include "Visitor/Visitor.h"
+
//#ifdef WIN32
// #include <Winsock2.h>
//#endif
@@ -1901,6 +1903,18 @@ void Parameter::setADKernel(std::string adKernel)
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+//add-methods
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+void Parameter::addActuator(Visitor* actuator)
+{
+ actuators.push_back(actuator);
+}
+void Parameter::addProbe(Visitor* probe)
+{
+ probes.push_back(probe);
+}
+
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
//get-methods
@@ -2303,6 +2317,14 @@ TempPressforBoundaryConditions* Parameter::getTempPressD()
{
return this->TempPressD;
}
+std::vector<Visitor*> Parameter::getActuators()
+{
+ return actuators;
+}
+std::vector<Visitor*> Parameter::getProbes()
+{
+ return probes;
+}
//unsigned int Parameter::getkInflowQ()
//{
// return this->kInflowQ;
@@ -2853,6 +2875,7 @@ std::string Parameter::getADKernel()
{
return adKernel;
}
+
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
void Parameter::setInitialCondition(std::function<void(real,real,real,real&,real&,real&,real&)> initialCondition)
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index ccb13e83b..486ef71cd 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -13,6 +13,7 @@
#include "LBM/LB.h"
#include "LBM/D3Q27.h"
#include "Calculation/PorousMedia.h"
+#include "Visitor/Visitor.h"
//#include "Output/LogWriter.hpp"
#include <cuda_runtime.h>
@@ -512,6 +513,12 @@ public:
void setADKernel(std::string adKernel);
+
+ //adder
+
+ void addActuator(Visitor* actuator);
+ void addProbe(Visitor* probes);
+
//getter
double* getForcesDouble();
real* getForcesHost();
@@ -668,6 +675,8 @@ public:
TempVelforBoundaryConditions* getTempVelD();
TempPressforBoundaryConditions* getTempPressH();
TempPressforBoundaryConditions* getTempPressD();
+ std::vector<Visitor*> getActuators();
+ std::vector<Visitor*> getProbes();
unsigned int getTimeDoCheckPoint();
unsigned int getTimeDoRestart();
bool getDoCheckPoint();
@@ -726,6 +735,7 @@ public:
std::string getADKernel();
+
~Parameter();
public:
@@ -791,6 +801,9 @@ private:
//Temperature Pressure
TempPressforBoundaryConditions *TempPressH, *TempPressD;
+ std::vector<Visitor*> actuators;
+ std::vector<Visitor*> probes;
+
//Drehung///////////////
real Phi, angularVelocity;
unsigned int startTurn;
@@ -823,6 +836,8 @@ private:
std::vector<std::string> possNeighborFilesSendX, possNeighborFilesSendY, possNeighborFilesSendZ;
std::vector<std::string> possNeighborFilesRecvX, possNeighborFilesRecvY, possNeighborFilesRecvZ;
bool isNeigborX, isNeigborY, isNeigborZ;
+
+
////////////////////////////////////////////////////////////////////////////
// initial condition
diff --git a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu
new file mode 100644
index 000000000..3dcbea259
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.cu
@@ -0,0 +1,177 @@
+#include "ActuatorLine.h"
+
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <helper_cuda.h>
+
+__global__ void interpolateVelocities(real* coordsX, real* coordsY, real* coordsZ, int* indices, int numberOfIndices);
+
+void ActuatorLine::init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaManager)
+{
+ this->allocBladeCoords(cudaManager);
+ this->initBladeCoords();
+ this->initBoundingSphere(para, cudaManager);
+}
+
+void ActuatorLine::visit(Parameter* para, int level, unsigned int t)
+{
+ if (level != this->level) return;
+
+ this->copyBladeCoordsHtoD();
+
+ unsigned int numberOfThreads = 128;
+ int Grid = (this->numberOfIndices/ numberOfThreads)+1;
+ int Grid1, Grid2;
+ if (Grid>512)
+ {
+ Grid1 = 512;
+ Grid2 = (Grid/Grid1)+1;
+ }
+ else
+ {
+ Grid1 = 1;
+ Grid2 = Grid;
+ }
+ dim3 grid(Grid1, Grid2);
+ dim3 threads(numberOfThreads, 1, 1 );
+
+ interpolateVelocities<<< grid, threads >>>(
+ para->getParD(this->level)->coordX_SP,
+ para->getParD(this->level)->coordY_SP,
+ para->getParD(this->level)->coordZ_SP,
+ this->boundingSphereIndicesD,
+ this->numberOfIndices);
+
+}
+
+void ActuatorLine::allocBladeCoords(CudaMemoryManager* cudaManager)
+{
+ checkCudaErrors( cudaMallocHost((void**) &this->bladeCoordsXH, this->mem_size_blades) );
+ checkCudaErrors( cudaMallocHost((void**) &this->bladeCoordsYH, this->mem_size_blades) );
+ checkCudaErrors( cudaMallocHost((void**) &this->bladeCoordsZH, this->mem_size_blades) );
+
+ checkCudaErrors( cudaMalloc((void**) &this->bladeCoordsXD, this->mem_size_blades) );
+ checkCudaErrors( cudaMalloc((void**) &this->bladeCoordsYD, this->mem_size_blades) );
+ checkCudaErrors( cudaMalloc((void**) &this->bladeCoordsZD, this->mem_size_blades) );
+
+ cudaManager->setMemsizeGPU(3.*this->mem_size_blades, false);
+}
+
+void ActuatorLine::initBladeCoords()
+{
+ real dx = 0.5f*this->diameter/this->nBladeNodes;
+ for(unsigned int node=0; node<this->nBladeNodes; node++)
+ {
+ this->bladeCoordsXH[node] = this->turbinePosX;
+ this->bladeCoordsYH[node] = this->turbinePosY;
+ this->bladeCoordsYH[node] = this->turbinePosZ+dx*node;
+
+ this->bladeCoordsXH[node+this->nBladeNodes] = this->turbinePosX;
+ this->bladeCoordsYH[node+this->nBladeNodes] = this->turbinePosY*0.5f*sqrt(3.0f)*dx*node;
+ this->bladeCoordsYH[node+this->nBladeNodes] = this->turbinePosZ-0.5f*dx*node;
+
+ this->bladeCoordsXH[node+2*this->nBladeNodes] = this->turbinePosX;
+ this->bladeCoordsYH[node+2*this->nBladeNodes] = this->turbinePosY-0.5f*sqrt(3.0f)*dx*node;
+ this->bladeCoordsYH[node+2*this->nBladeNodes] = this->turbinePosZ-0.5f*dx*node;
+ }
+}
+
+void ActuatorLine::copyBladeCoordsHtoD()
+{
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsXD, this->bladeCoordsXH, this->mem_size_blades, cudaMemcpyHostToDevice) );
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsYD, this->bladeCoordsYH, this->mem_size_blades, cudaMemcpyHostToDevice) );
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsZD, this->bladeCoordsZH, this->mem_size_blades, cudaMemcpyHostToDevice) );
+}
+
+void ActuatorLine::copyBladeCoordsDtoH()
+{
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsXH, this->bladeCoordsXD, this->mem_size_blades, cudaMemcpyDeviceToHost) );
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsYH, this->bladeCoordsYD, this->mem_size_blades, cudaMemcpyDeviceToHost) );
+ checkCudaErrors( cudaMemcpy(this->bladeCoordsZH, this->bladeCoordsZD, this->mem_size_blades, cudaMemcpyDeviceToHost) );
+}
+
+void ActuatorLine::freeBladeCoords()
+{
+ checkCudaErrors( cudaFree(this->bladeCoordsXD) );
+ checkCudaErrors( cudaFree(this->bladeCoordsYD) );
+ checkCudaErrors( cudaFree(this->bladeCoordsZD) );
+
+ checkCudaErrors( cudaFreeHost(this->bladeCoordsXH) );
+ checkCudaErrors( cudaFreeHost(this->bladeCoordsYH) );
+ checkCudaErrors( cudaFreeHost(this->bladeCoordsZH) );
+}
+void ActuatorLine::rotateBlades(real angle)
+{
+ for(unsigned int node=0; node<this->nBladeNodes*this->nBlades; node++)
+ {
+ this->bladeCoordsYH[node] = this->bladeCoordsYH[node]*cos(angle)-this->bladeCoordsZH[node]*sin(angle);
+ this->bladeCoordsZH[node] = this->bladeCoordsYH[node]*sin(angle)+this->bladeCoordsZH[node]*cos(angle);
+ }
+}
+
+void ActuatorLine::initBoundingSphere(Parameter* para, CudaMemoryManager* cudaManager)
+{
+ // Actuator line exists only on 1 level
+ std::vector<int> nodesInSphere;
+
+ for (int 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];
+ const real coordZ = para->getParH(this->level)->coordZ_SP[j];
+ const real dist = sqrt(pow(coordX-this->turbinePosX,2)+pow(coordY-this->turbinePosY,2)+pow(coordZ-this->turbinePosZ,2));
+ if(dist < 0.6*this->diameter)
+ {
+ printf("indx in: %i \n", j);
+ nodesInSphere.push_back(j);
+ }
+ }
+
+
+ this->numberOfIndices = nodesInSphere.size();
+ this->mem_size_boundingSphere = sizeof(int)*this->numberOfIndices;
+ this->allocSphereIndices(cudaManager);
+ this->boundingSphereIndicesH = nodesInSphere.data();
+ this->copySphereIndices();
+
+}
+
+void ActuatorLine::allocSphereIndices(CudaMemoryManager* cudaManager)
+{
+ printf("mem size sphere %i", (int)this->mem_size_boundingSphere);
+
+ checkCudaErrors( cudaMallocHost((void**) &(this->boundingSphereIndicesH), this->mem_size_boundingSphere));
+ checkCudaErrors( cudaMalloc((void**) &(this->boundingSphereIndicesD), this->mem_size_boundingSphere));
+ cudaManager->setMemsizeGPU(this->mem_size_boundingSphere, false);
+}
+
+void ActuatorLine::copySphereIndices()
+{
+ checkCudaErrors( cudaMemcpy(this->boundingSphereIndicesD, this->boundingSphereIndicesH, this->mem_size_boundingSphere, cudaMemcpyHostToDevice) );
+}
+
+void ActuatorLine::freeSphereIndices()
+{
+ checkCudaErrors( cudaFreeHost(this->boundingSphereIndicesH) );
+ checkCudaErrors( cudaFree(this->boundingSphereIndicesD) );
+}
+
+__global__ void interpolateVelocities(real* coordsX, real* coordsY, real* coordsZ, int* indices, int numberOfIndices)
+{
+ const uint x = threadIdx.x; // Globaler x-Index
+ const uint y = blockIdx.x; // Globaler y-Index
+ const uint z = blockIdx.y; // Globaler z-Index
+
+ const uint nx = blockDim.x;
+ const uint ny = gridDim.x;
+
+ const uint index = nx*(ny*z + y) + x;
+
+ if(index>=numberOfIndices) return;
+
+ if(index==0||index+1==numberOfIndices)
+ {
+ printf("idx, x, y, z: %i, %f, %f, %f \n", indices[index], coordsX[index],coordsY[index],coordsZ[index]);
+ }
+
+}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h
new file mode 100644
index 000000000..0a759a5f8
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/Visitor/ActuatorLine.h
@@ -0,0 +1,78 @@
+#ifndef ActuatorLine_H
+#define ActuatorLine_H
+
+#include "Visitor.h"
+#include "Parameter/Parameter.h"
+#include "PointerDefinitions.h"
+#include "GridGenerator/grid/GridBuilder/GridBuilder.h"
+
+class ActuatorLine : public Visitor
+{
+public:
+ ActuatorLine(
+ const unsigned int &_nBlades,
+ const real &_density,
+ const unsigned int &_nBladeNodes,
+ const real &_epsilon,
+ real &_turbinePosX, real &_turbinePosY, real &_turbinePosZ,
+ const real &_diameter,
+ int &_level
+ ) : nBlades(_nBlades),
+ density(_density),
+ nBladeNodes(_nBladeNodes),
+ epsilon(_epsilon),
+ turbinePosX(_turbinePosX), turbinePosY(_turbinePosY), turbinePosZ(_turbinePosZ),
+ diameter(_diameter),
+ level(_level),
+ Visitor()
+ {
+ this->mem_size_blades = sizeof(real)*this->nBladeNodes*this->nBlades;
+ this->omega = 0.0f;
+ this->azimuth = 0.0f;
+ }
+
+ virtual ~ActuatorLine()
+ {
+ this->freeBladeCoords();
+ this->freeSphereIndices();
+ }
+
+ void visit(Parameter* para, int level, unsigned int t);
+ void initBoundingSphere(Parameter* para, CudaMemoryManager* cudaManager);
+ void init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaManager);
+
+private:
+
+ void allocBladeCoords(CudaMemoryManager* cudaManager);
+ void initBladeCoords();
+ void copyBladeCoordsHtoD();
+ void copyBladeCoordsDtoH();
+ void freeBladeCoords();
+ void rotateBlades(real angle);
+
+ void allocSphereIndices(CudaMemoryManager* cudaManager);
+ void copySphereIndices();
+ void freeSphereIndices();
+
+private:
+ const real density;
+ real turbinePosX, turbinePosY, turbinePosZ;
+ real* bladeCoordsXH, * bladeCoordsYH, * bladeCoordsZH;
+ real* bladeCoordsXD, * bladeCoordsYD, * bladeCoordsZD;
+ int* boundingSphereIndicesH, * boundingSphereIndicesD;
+ real omega, azimuth;
+ const real diameter;
+ const unsigned int nBladeNodes;
+ const unsigned int nBlades;
+ const real epsilon; // in m
+ const int level;
+ int numberOfIndices;
+ size_t mem_size_blades;
+ size_t mem_size_boundingSphere;
+
+ unsigned int* indicesBoundingBox;
+
+
+};
+
+#endif
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/Visitor/Visitor.h b/src/gpu/VirtualFluids_GPU/Visitor/Visitor.h
new file mode 100644
index 000000000..40c3dafa3
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/Visitor/Visitor.h
@@ -0,0 +1,58 @@
+#ifndef Visitor_H
+#define Visitor_H
+
+#include <vector>
+#include <string>
+
+#include "PointerDefinitions.h"
+#include "Core/DataTypes.h"
+#include "VirtualFluids_GPU_export.h"
+#include "DataStructureInitializer/GridProvider.h"
+#include "GPU/CudaMemoryManager.h"
+
+#include <cassert>
+
+class Parameter;
+
+
+class VIRTUALFLUIDS_GPU_EXPORT Visitor
+{
+private:
+ SPtr<Parameter> para;
+
+
+protected:
+ Visitor()
+ {
+ this->updateInterval = 1;
+ }
+
+public:
+ virtual ~Visitor() {}
+
+ virtual void visit(Parameter* para, int level, unsigned int t)=0;
+ virtual void init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaManager)=0;
+
+ ////////////////////////////////////////////////////////////////////////////
+ /// \brief setUpdateInterval
+ /// \param _updateInterval
+ /// \note usage: setUpdateInterval(this->tout);
+ /// \note usage: setUpdateInterval(div_ru(this->tout,10U));
+ ///
+ void setUpdateInterval(const uint &_updateInterval)
+ {
+ assert(_updateInterval>0);
+ this->updateInterval = _updateInterval;
+ }
+
+ bool isDue(const uint &tLB) const
+ {
+ return (tLB%this->updateInterval==0);
+ }
+
+protected:
+ uint updateInterval; ///< update interval in number of timesteps of the coars patch (1 = each time step)
+
+};
+
+#endif
\ No newline at end of file
--
GitLab