diff --git a/CMake/cmake_config_files/ARAGORN.config.cmake b/CMake/cmake_config_files/ARAGORN.config.cmake
index 7fc74bbcbf39890ebc26d328b4916575c8471a33..5af2b3809558ae9caab753e43f6082c753607086 100644
--- a/CMake/cmake_config_files/ARAGORN.config.cmake
+++ b/CMake/cmake_config_files/ARAGORN.config.cmake
@@ -9,8 +9,8 @@ set(CMAKE_CUDA_ARCHITECTURES 86) # Nvidia GeForce RTX 3060
# add invidual apps here
set(GPU_APP "apps/gpu/")
list(APPEND USER_APPS
- "${GPU_APP}DrivenCavityMultiGPU"
- "${GPU_APP}SphereScaling"
+ # "${GPU_APP}DrivenCavityMultiGPU"
+ # "${GPU_APP}SphereScaling"
# "${GPU_APP}MusselOyster"
"${GPU_APP}RotatingGrid"
)
diff --git a/apps/gpu/RotatingGrid/RotatingGrid.cpp b/apps/gpu/RotatingGrid/RotatingGrid.cpp
index 3310f13e5aefe564955732538046369b55a7866a..59d0cd83ec1e33d2b60f1b569a6b41c50f3590ef 100644
--- a/apps/gpu/RotatingGrid/RotatingGrid.cpp
+++ b/apps/gpu/RotatingGrid/RotatingGrid.cpp
@@ -92,7 +92,7 @@ int main()
const uint nx = 64;
const uint timeStepOut = 1;
- const uint timeStepEnd = 1;
+ const uint timeStepEnd = 10;
//////////////////////////////////////////////////////////////////////////
// compute parameters in lattice units
@@ -110,7 +110,7 @@ int main()
gridBuilder->addCoarseGrid(-1.5 * L, -0.5 * L, -0.5 * L, 1.5 * L, 0.5 * L, 0.5 * L, dx);
- if (rotOrInt == Rot) gridBuilder->addGridRotatingGrid(std::make_shared<Cylinder>(0.0, 0.0, 0.0, 0.3 * L, 1. * L, Axis::x));
+ if (rotOrInt == Rot) gridBuilder->addGridRotatingGrid(std::make_shared<Cylinder>(0.1, 0.1, 0.1, 0.25 * L, 1. * L, Axis::x));
if (rotOrInt == Int) gridBuilder->addGrid(std::make_shared<Cylinder>(0.0, 0.0, 0.0, 0.3 * L, 1. * L, Axis::x), 1);
GridScalingFactory scalingFactory = GridScalingFactory();
@@ -205,7 +205,6 @@ int main()
// NeighborDebugWriter::writeNeighborLinkLinesDebug(para.get());
SPtr<ParameterRotatingGrid> paraRot = para->getRotatingGridParameter();
- paraRot->transformNestedToBase({para->getParH(1)->coordinateX, para->getParH(1)->coordinateY, para->getParH(1)->coordinateZ});
sim.run();
} catch (const spdlog::spdlog_ex &ex) {
diff --git a/src/gpu/GridGenerator/grid/GridImp.cpp b/src/gpu/GridGenerator/grid/GridImp.cpp
index 62e2c95c1811a76487a490c81ca028288bc0a4bc..ab5a3ac21efd68f82f9fe345cc7f41637c838340 100644
--- a/src/gpu/GridGenerator/grid/GridImp.cpp
+++ b/src/gpu/GridGenerator/grid/GridImp.cpp
@@ -1165,15 +1165,15 @@ void GridImp::findGridInterfaceForRotatingGrid(SPtr<Grid> rotatingGrid)
for (uint indexOnStaticGrid = 0; indexOnStaticGrid < this->getSize(); indexOnStaticGrid++)
this->findOverlapStopper(indexOnStaticGrid, *rotatingGridImp);
- // // change interface gap nodes to fluid nodes on static grid
- // for (uint indexOnStaticGrid = 0; indexOnStaticGrid < this->getSize(); indexOnStaticGrid++)
- // if (this->getField().isInterpolationGapNode(indexOnStaticGrid))
- // this->getField().setFieldEntry(indexOnStaticGrid, FLUID);
+ // change interface gap nodes to fluid nodes on static grid
+ for (uint indexOnStaticGrid = 0; indexOnStaticGrid < this->getSize(); indexOnStaticGrid++)
+ if (this->getField().isInterpolationGapNode(indexOnStaticGrid))
+ this->getField().setFieldEntry(indexOnStaticGrid, FLUID);
- // // change interface gap nodes to fluid nodes on rotating grid
- // for (uint indexOnRotatingGrid = 0; indexOnRotatingGrid < rotatingGridImp->getSize(); indexOnRotatingGrid++)
- // if (rotatingGridImp->getField().isInterpolationGapNode(indexOnRotatingGrid))
- // rotatingGridImp->getField().setFieldEntry(indexOnRotatingGrid, FLUID);
+ // change interface gap nodes to fluid nodes on rotating grid
+ for (uint indexOnRotatingGrid = 0; indexOnRotatingGrid < rotatingGridImp->getSize(); indexOnRotatingGrid++)
+ if (rotatingGridImp->getField().isInterpolationGapNode(indexOnRotatingGrid))
+ rotatingGridImp->getField().setFieldEntry(indexOnRotatingGrid, FLUID);
VF_LOG_TRACE(" ... done.");
}
@@ -1185,6 +1185,8 @@ void GridImp::repairGridInterfaceOnMultiGPU(SPtr<Grid> fineGrid)
void GridImp::limitToSubDomain(SPtr<BoundingBox> subDomainBox)
{
+ VF_LOG_TRACE("Limit to subdomain box.");
+
for( uint index = 0; index < this->size; index++ ){
real x, y, z;
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index cf1aaa3988e0809dd0995a18139ab7dcef75989f..9e36a27f2aac54d669327298a24e6b8e4244be61 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -9,6 +9,7 @@
#include "KernelManager/GridScalingKernelManager.h"
#include "TurbulenceModels/TurbulenceModelFactory.h"
#include "Kernel/Kernel.h"
+#include "Parameter/ParameterRotatingGrid.h"
#include "CollisionStrategy.h"
#include "RefinementStrategy.h"
@@ -52,9 +53,21 @@ void UpdateGrid27::updateGrid(int level, unsigned int t)
this->preCollisionBC(level, t);
//////////////////////////////////////////////////////////////////////////
- if( level != para->getFine() )
- {
- refinement(this, para.get(), level);
+ if (level != para->getFine()) {
+ SPtr<ParameterRotatingGrid> paraRot = para->getRotatingGridParameter();
+ if ( paraRot != nullptr) {
+ // rotation
+ paraRot->parameterRotHost->gridAngle[0] += paraRot->parameterRotHost->angularVelocity[0];
+ paraRot->parameterRotHost->gridAngle[1] += paraRot->parameterRotHost->angularVelocity[1];
+ paraRot->parameterRotHost->gridAngle[2] += paraRot->parameterRotHost->angularVelocity[2];
+ paraRot->parameterRotDevice->gridAngle[0] += paraRot->parameterRotDevice->angularVelocity[0];
+ paraRot->parameterRotDevice->gridAngle[1] += paraRot->parameterRotDevice->angularVelocity[1];
+ paraRot->parameterRotDevice->gridAngle[2] += paraRot->parameterRotDevice->angularVelocity[2];
+ VF_LOG_DEBUG("gridAngleX = {}", paraRot->parameterRotDevice->gridAngle[0]);
+ rotationInterpolation(level);
+ } else {
+ refinement(this, para.get(), level);
+ }
}
//////////////////////////////////////////////////////////////////////////
@@ -63,6 +76,19 @@ void UpdateGrid27::updateGrid(int level, unsigned int t)
}
+void UpdateGrid27::rotationInterpolation(int level)
+{
+ VF_LOG_WARNING("Interpolation");
+ // base to nested
+ InterpolateStaticToRotating(
+ para->getParD(level).get(),
+ para->getParD(level+1).get(),
+ para->getRotatingGridParameter()->parameterRotDevice.get(),
+ ¶->getParD(level)->coarseToFine,
+ para->getParD(level)->neighborCoarseToFine);
+ // nested to base
+}
+
void UpdateGrid27::collisionAllNodes(int level, unsigned int t)
{
kernels.at(level)->run();
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index 9de7e73ec03a0f00542dd5b718ed6f210399a18a..5218fdf8081e98592673efb0b78bdaf9f1faa46b 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -59,6 +59,8 @@ private:
void interactWithActuators(int level, unsigned int t);
void interactWithProbes(int level, unsigned int t);
+ void rotationInterpolation(int level);
+
private:
CollisionStrategy collision;
friend class CollisionAndExchange_noStreams_indexKernel;
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index e96c96ec2f38c4d27f7c51177e4a3c21f37579e8..4336928e8b0dbeedcc864ac892d6e60ac0e3a785 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -149,6 +149,7 @@ void GridProvider::freeMemoryOnHost()
cudaMemoryManager->cudaFreeCoord(level);
cudaMemoryManager->cudaFreeSP(level);
}
+ cudaMemoryManager->cudaFreeCoordRotation(1);
}
void GridProvider::cudaCopyDataToHost(int level)
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index 19d9890e9fa124ad8125bbb4f136bdce3f456118..ee7aa332a90c77784f076f3ede8cc9a1472e8876 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -50,9 +50,9 @@ void GridGenerator::initalGridInformations()
builder->findFluidNodes(para->getUseStreams());
std::vector<int> gridX, gridY, gridZ;
std::vector<int> distX, distY, distZ;
- const int numberOfGridLevels = builder->getNumberOfGridLevels();
+ const uint numberOfGridLevels = builder->getNumberOfGridLevels();
builder->getGridInformations(gridX, gridY, gridZ, distX, distY, distZ);
- para->setMaxLevel(numberOfGridLevels);
+ para->setMaxLevel((int) numberOfGridLevels);
para->setGridX(gridX);
para->setGridY(gridY);
para->setGridZ(gridZ);
@@ -108,12 +108,23 @@ void GridGenerator::allocArrays_CoordNeighborGeo()
cudaMemoryManager->cudaCopyBodyForce(level);
}
+
if (builder->getUseRotatingGrid()) {
+ VF_LOG_INFO("Init parameters for rotating grid.");
SPtr<const Object> cylinder = builder->getGrid(1)->getObject();
const std::array<real, 3> centerPointOfCylinder = { (real)cylinder->getX1Centroid(), (real)cylinder->getX2Centroid(),
(real)cylinder->getX3Centroid() };
- auto parameterRotatingGrid = std::make_shared<ParameterRotatingGrid>(centerPointOfCylinder, Axis::x);
- para->setRotatingGridParameter(std::move(parameterRotatingGrid));
+ para->setRotatingGridParameter(
+ std::make_shared<ParameterRotatingGrid>(centerPointOfCylinder, Axis::x, para->getParH(1)->numberOfNodes));
+ cudaMemoryManager->cudaAllocCoordRotation(1);
+ para->fillCoordinateVectorsForRotatingGrid(1);
+ cudaMemoryManager->cudaCopyCoordRotation(1);
+ }
+
+ for (int i = 0; i <= para->getMaxLevel(); i++) {
+ para->getParH(i)->gridSpacing = builder->getGrid(i)->getDelta();
+ para->getParD(i)->gridSpacing = builder->getGrid(i)->getDelta();
+ VF_LOG_WARNING("{}", builder->getGrid(i)->getDelta());
}
VF_LOG_INFO("Number of Nodes: {}", numberOfNodesGlobal);
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 64943d19ff54bfaa174d0728a96c517f6605d565..c2dfc7512b21cf54ab7da7aaeb74ccc6c4563f00 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -5,6 +5,7 @@
#include <math.h>
#include <Parameter/Parameter.h>
+#include <Parameter/ParameterRotatingGrid.h>
#include "Parameter/CudaStreamManager.h"
#include "PreCollisionInteractor/ActuatorFarm.h"
@@ -46,31 +47,74 @@ void CudaMemoryManager::cudaCopyMedianPrint(int lev)
}
void CudaMemoryManager::cudaAllocCoord(int lev)
{
- //Host
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateX ), parameter->getParH(lev)->memSizeRealLBnodes ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateY ), parameter->getParH(lev)->memSizeRealLBnodes ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateZ ), parameter->getParH(lev)->memSizeRealLBnodes ));
- //Device (spinning ship + uppsala)
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateX ), parameter->getParH(lev)->memSizeRealLBnodes ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateY ), parameter->getParH(lev)->memSizeRealLBnodes ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateZ ), parameter->getParH(lev)->memSizeRealLBnodes ));
- //////////////////////////////////////////////////////////////////////////
- double tmp = 3. * (double)parameter->getParH(lev)->memSizeRealLBnodes;
- setMemsizeGPU(tmp, false);
+ //Host
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateX ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateY ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->coordinateZ ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ //Device (spinning ship + uppsala)
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateX ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateY ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->coordinateZ ), parameter->getParH(lev)->memSizeRealLBnodes ));
+ //////////////////////////////////////////////////////////////////////////
+ double tmp = 3. * (double)parameter->getParH(lev)->memSizeRealLBnodes;
+ setMemsizeGPU(tmp, false);
}
void CudaMemoryManager::cudaCopyCoord(int lev)
{
- //copy host to device
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateX, parameter->getParH(lev)->coordinateX, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateY, parameter->getParH(lev)->coordinateY, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateZ, parameter->getParH(lev)->coordinateZ, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
+ //copy host to device
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateX, parameter->getParH(lev)->coordinateX, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateY, parameter->getParH(lev)->coordinateY, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->coordinateZ, parameter->getParH(lev)->coordinateZ, parameter->getParH(lev)->memSizeRealLBnodes , cudaMemcpyHostToDevice));
+}
+void CudaMemoryManager::cudaCopyCoordDeviceToHost(int lev)
+{
+ //copy device to host
+ checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->coordinateX, parameter->getParD(lev)->coordinateX, parameter->getParH(lev)->memSizeRealLBnodes, cudaMemcpyDeviceToHost));
+ checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->coordinateY, parameter->getParD(lev)->coordinateY, parameter->getParH(lev)->memSizeRealLBnodes, cudaMemcpyDeviceToHost));
+ checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->coordinateZ, parameter->getParD(lev)->coordinateZ, parameter->getParH(lev)->memSizeRealLBnodes, cudaMemcpyDeviceToHost));
}
void CudaMemoryManager::cudaFreeCoord(int lev)
{
- checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateX ));
- checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateY ));
- checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateZ ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateX ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateY ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coordinateZ ));
}
+void CudaMemoryManager::cudaAllocCoordRotation(int lev)
+{
+//Host
+ unsigned long long memSize = parameter->getRotatingGridParameter()->parameterRotHost->memorySizeOfNestedCoordinates;
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesX), memSize ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesY), memSize ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesZ), memSize ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesX), memSize ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesY), memSize ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesZ), memSize ));
+ //////////////////////////////////////////////////////////////////////////
+ setMemsizeGPU(3. * memSize, false);
+}
+void CudaMemoryManager::cudaCopyCoordRotation(int lev)
+{
+ //copy host to device
+ unsigned long long memSize = parameter->getRotatingGridParameter()->parameterRotHost->memorySizeOfNestedCoordinates;
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesX, parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesX, memSize , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesY, parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesY, memSize , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesZ, parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesZ, memSize , cudaMemcpyHostToDevice));
+}
+void CudaMemoryManager::cudaCopyCoordRotationDeviceToHost(int lev)
+{
+ //copy device to host
+ unsigned long long memSize = parameter->getRotatingGridParameter()->parameterRotHost->memorySizeOfNestedCoordinates;
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesX, parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesX, memSize, cudaMemcpyDeviceToHost));
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesY, parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesY, memSize, cudaMemcpyDeviceToHost));
+ checkCudaErrors( cudaMemcpy(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesZ, parameter->getRotatingGridParameter()->parameterRotDevice->nestedCoordinatesZ, memSize, cudaMemcpyDeviceToHost));
+}
+void CudaMemoryManager::cudaFreeCoordRotation(int lev)
+{
+ checkCudaErrors( cudaFreeHost(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesX ));
+ checkCudaErrors( cudaFreeHost(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesY ));
+ checkCudaErrors( cudaFreeHost(parameter->getRotatingGridParameter()->parameterRotHost->nestedCoordinatesZ ));
+}
+
void CudaMemoryManager::cudaAllocBodyForce(int lev)
{
//Host
@@ -84,7 +128,6 @@ void CudaMemoryManager::cudaAllocBodyForce(int lev)
//////////////////////////////////////////////////////////////////////////
double tmp = 3. * (double)parameter->getParH(lev)->memSizeRealLBnodes;
setMemsizeGPU(tmp, false);
-
}
void CudaMemoryManager::cudaCopyBodyForce(int lev)
{
@@ -1124,6 +1167,13 @@ void CudaMemoryManager::cudaCopyInterfaceCF(int lev)
checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->coarseToFine.coarseCellIndices, parameter->getParH(lev)->coarseToFine.coarseCellIndices, mem_size_kCF, cudaMemcpyHostToDevice));
checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->coarseToFine.fineCellIndices, parameter->getParH(lev)->coarseToFine.fineCellIndices, mem_size_kCF, cudaMemcpyHostToDevice));
}
+void CudaMemoryManager::cudaCopyInterfaceCFDeviceToHost(int lev)
+{
+ uint mem_size_kCF = sizeof(uint) * parameter->getParH(lev)->coarseToFine.numberOfCells;
+
+ checkCudaErrors(cudaMemcpy(parameter->getParH(lev)->coarseToFine.coarseCellIndices, parameter->getParD(lev)->coarseToFine.coarseCellIndices, mem_size_kCF, cudaMemcpyDeviceToHost));
+ checkCudaErrors(cudaMemcpy(parameter->getParH(lev)->coarseToFine.fineCellIndices , parameter->getParD(lev)->coarseToFine.fineCellIndices, mem_size_kCF, cudaMemcpyDeviceToHost));
+}
void CudaMemoryManager::cudaFreeInterfaceCF(int lev)
{
checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->coarseToFine.coarseCellIndices));
@@ -1151,6 +1201,12 @@ void CudaMemoryManager::cudaCopyInterfaceFC(int lev)
checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->fineToCoarse.fineCellIndices, parameter->getParH(lev)->fineToCoarse.fineCellIndices, mem_size_kFC, cudaMemcpyHostToDevice));
checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->fineToCoarse.coarseCellIndices, parameter->getParH(lev)->fineToCoarse.coarseCellIndices, mem_size_kFC, cudaMemcpyHostToDevice));
}
+void CudaMemoryManager::cudaCopyInterfaceFCDeviceToHost(int lev){
+ uint mem_size_kFC = sizeof(uint) * parameter->getParH(lev)->fineToCoarse.numberOfCells;
+
+ checkCudaErrors(cudaMemcpy(parameter->getParH(lev)->fineToCoarse.fineCellIndices , parameter->getParD(lev)->fineToCoarse.fineCellIndices , mem_size_kFC, cudaMemcpyDeviceToHost));
+ checkCudaErrors(cudaMemcpy(parameter->getParH(lev)->fineToCoarse.coarseCellIndices, parameter->getParD(lev)->fineToCoarse.coarseCellIndices, mem_size_kFC, cudaMemcpyDeviceToHost));
+}
void CudaMemoryManager::cudaCheckInterfaceFCBulk(int lev)
{
// only use for testing!
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
index 1d01c227e780ec564b9eb1506f1a0183769d64b9..20c5a1145c6002803d54302e976cc22ddf4e3e10 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.h
@@ -49,6 +49,12 @@ public:
void cudaAllocCoord(int lev);
void cudaCopyCoord(int lev);
void cudaFreeCoord(int lev);
+ void cudaCopyCoordDeviceToHost(int lev);
+
+ void cudaAllocCoordRotation(int lev);
+ void cudaCopyCoordRotation(int lev);
+ void cudaFreeCoordRotation(int lev);
+ void cudaCopyCoordRotationDeviceToHost(int lev);
void cudaAllocBodyForce(int lev);
void cudaCopyBodyForce(int lev);
@@ -164,10 +170,12 @@ public:
void cudaAllocInterfaceCF(int lev);
void cudaCopyInterfaceCF(int lev);
+ void cudaCopyInterfaceCFDeviceToHost(int lev);
void cudaFreeInterfaceCF(int lev);
void cudaAllocInterfaceFC(int lev);
void cudaCopyInterfaceFC(int lev);
+ void cudaCopyInterfaceFCDeviceToHost(int lev);
void cudaCheckInterfaceFCBulk(int lev);
void cudaFreeInterfaceFC(int lev);
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
index b449bcca404b54469b52f7b383f0615632c4d06d..f33013fe9bb6777522c388c6158d0929b37cfc5c 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Interface.h
@@ -25,6 +25,7 @@
#endif
struct LBMSimulationParameter;
+struct ParameterRotatingGridSimulation;
class Parameter;
//////////////////////////////////////////////////////////////////////////
@@ -1622,6 +1623,17 @@ void ScaleCF_RhoSq_comp_27(LBMSimulationParameter * parameterDeviceC, LBMSimulat
template<bool hasTurbulentViscosity> void ScaleCF_compressible(LBMSimulationParameter * parameterDeviceC, LBMSimulationParameter* parameterDeviceF, ICells * interpolationCellsCoarseToFine, ICellNeigh &neighborCoarseToFine, CUstream_st *stream);
+void InterpolateStaticToRotating(
+ LBMSimulationParameter *parameterDeviceS,
+ LBMSimulationParameter *parameterDeviceR,
+ ParameterRotatingGridSimulation *paraRotDevice,
+ ICells *baseToNested,
+ ICellNeigh &neighborBaseToNested);
+
+void UpdateGlobalCoordinates(
+ LBMSimulationParameter *parameterDeviceR,
+ ParameterRotatingGridSimulation *paraRotDevice
+);
void ScaleCF_RhoSq_3rdMom_comp_27( real* DC,
real* DF,
unsigned int* neighborCX,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
index 74ebf3bea73c221207d3dda7a6a2f29de083ffde..5b43dabe4db102e3e47e3fe07699835b9f58d9aa 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
+++ b/src/gpu/VirtualFluids_GPU/GPU/GPU_Kernels.cuh
@@ -1868,6 +1868,46 @@ template<bool hasTurbulentViscosity> __global__ void scaleCF_compressible(
real* turbulentViscosityFine,
ICellNeigh offsetCF);
+__global__ void interpolateStaticToRotating(
+ unsigned int numberOfInterfaceNodes,
+ unsigned int *indicesStaticMMM,
+ unsigned int *indicesRotating,
+ const real *coordDestinationX,
+ const real *coordDestinationY,
+ const real *coordDestinationZ,
+ const real *coordSourceX,
+ const real *coordSourceY,
+ const real *coordSourceZ,
+ const uint *neighborXstatic,
+ const uint *neighborYstatic,
+ const uint *neighborZstatic,
+ const uint *neighborMMMstatic,
+ real centerCoordX,
+ real centerCoordY,
+ real centerCoordZ,
+ real angleX,
+ real angleY,
+ real angleZ,
+ real angularVelocityX,
+ real angularVelocityY,
+ real angularVelocityZ,
+ real dx);
+
+__global__ void updateGlobalCoordinates(
+ unsigned int numberOfNodes,
+ real *globalX,
+ real *globalY,
+ real *globalZ,
+ const real *localX,
+ const real *localY,
+ const real *localZ,
+ real centerCoordX,
+ real centerCoordY,
+ real centerCoordZ,
+ real angleX,
+ real angleY,
+ real angleZ);
+
__global__ void scaleCF_RhoSq_3rdMom_comp_27(real* DC,
real* DF,
unsigned int* neighborCX,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticRotating.cu b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticRotating.cu
new file mode 100644
index 0000000000000000000000000000000000000000..2714f2324b342cb0171d8940011b364d59a6c984
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticRotating.cu
@@ -0,0 +1,129 @@
+//=======================================================================================
+// ____ ____ __ ______ __________ __ __ __ __
+// \ \ | | | | | _ \ |___ ___| | | | | / \ | |
+// \ \ | | | | | |_) | | | | | | | / \ | |
+// \ \ | | | | | _ / | | | | | | / /\ \ | |
+// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____
+// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______|
+// \ \ | | ________________________________________________________________
+// \ \ | | | ______________________________________________________________|
+// \ \| | | | __ __ __ __ ______ _______
+// \ | | |_____ | | | | | | | | | _ \ / _____)
+// \ | | _____| | | | | | | | | | | \ \ \_______
+// \ | | | | |_____ | \_/ | | | | |_/ / _____ |
+// \ _____| |__| |________| \_______/ |__| |______/ (_______/
+//
+// This file is part of VirtualFluids. VirtualFluids is free software: you can
+// redistribute it and/or modify it under the terms of the GNU General Public
+// License as published by the Free Software Foundation, either version 3 of
+// the License, or (at your option) any later version.
+//
+// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT
+// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+// for more details.
+//
+// You should have received a copy of the GNU General Public License along
+// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
+//
+//! \file scaleCF_compressible.cu
+//! \ingroup GPU/GridScaling
+//! \author Martin Schoenherr, Anna Wellmann
+//=======================================================================================
+
+#include "DataTypes.h"
+#include "LBM/GPUHelperFunctions/KernelUtilities.h"
+#include "LBM/GPUHelperFunctions/ChimeraTransformation.h"
+#include "LBM/GPUHelperFunctions/ScalingUtilities.h"
+#include "LBM/GPUHelperFunctions/GridTraversion.h"
+#include "LBM/GPUHelperFunctions/CoordinateTransformation.h"
+
+#include <lbm/refinement/Coefficients.h>
+#include <lbm/refinement/InterpolationCF.h>
+
+__global__ void interpolateStaticToRotating(
+ unsigned int numberOfInterfaceNodes,
+ unsigned int * indicesStaticCell,
+ unsigned int *indicesRotating,
+ const real *coordDestinationX,
+ const real *coordDestinationY,
+ const real *coordDestinationZ,
+ const real *coordSourceX,
+ const real *coordSourceY,
+ const real *coordSourceZ,
+ const uint *neighborXstatic,
+ const uint *neighborYstatic,
+ const uint *neighborZstatic,
+ const uint *neighborMMMstatic,
+ real centerCoordX,
+ real centerCoordY,
+ real centerCoordZ,
+ real angleX,
+ real angleY,
+ real angleZ,
+ real angularVelocityX,
+ real angularVelocityY,
+ real angularVelocityZ,
+ real dx)
+{
+ const unsigned listIndex = vf::gpu::getNodeIndex();
+
+ if (listIndex >= numberOfInterfaceNodes) return;
+
+ uint destinationIndex = indicesRotating[listIndex];
+ uint sourceIndex = indicesStaticCell[listIndex];
+ uint indexNeighborMMMsource = neighborMMMstatic[sourceIndex];
+
+ real globalX;
+ real globalY;
+ real globalZ;
+
+ transformRotatingToGlobal(globalX, globalY, globalZ, coordDestinationX[destinationIndex],
+ coordDestinationY[destinationIndex], coordDestinationZ[destinationIndex], centerCoordX,
+ centerCoordY, centerCoordZ, angleX, angleY, angleZ);
+
+ // printf("coordDestinationX %.4f, coordSourceX %.4f, d-s %.4f \n", globalX, coordSourceX[sourceIndex],
+ // globalX - coordSourceX[sourceIndex]);
+
+ indicesStaticCell[listIndex] = traverseSourceCell(
+ globalX, globalY, globalZ,
+ indexNeighborMMMsource, coordSourceX[indexNeighborMMMsource], coordSourceY[indexNeighborMMMsource],
+ coordSourceZ[indexNeighborMMMsource], neighborXstatic, neighborYstatic, neighborZstatic, dx);
+
+ // if (sourceIndex != indicesStaticCell[listIndex]) {
+ // printf("sourceIndex: old = %d, new = %d\t", indicesStaticCell[listIndex], sourceIndex);
+ // }
+}
+
+__global__ void updateGlobalCoordinates(
+ unsigned int numberOfNodes,
+ real *globalX,
+ real *globalY,
+ real *globalZ,
+ const real *localX,
+ const real *localY,
+ const real *localZ,
+ real centerCoordX,
+ real centerCoordY,
+ real centerCoordZ,
+ real angleX,
+ real angleY,
+ real angleZ)
+{
+ const unsigned nodeIndex = vf::gpu::getNodeIndex();
+
+ if (nodeIndex >= numberOfNodes) return;
+
+ real globalXtemp;
+ real globalYtemp;
+ real globalZtemp;
+
+ transformRotatingToGlobal(globalXtemp, globalYtemp, globalZtemp, localX[nodeIndex], localY[nodeIndex], localZ[nodeIndex], centerCoordX,
+ centerCoordY, centerCoordZ, angleX, angleY, angleZ);
+
+ // printf("%.3f\t", localY[nodeIndex]);
+
+ globalX[nodeIndex] = globalXtemp;
+ globalY[nodeIndex] = globalYtemp;
+ globalZ[nodeIndex] = globalZtemp;
+}
\ No newline at end of file
diff --git a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
index e9bd3ea551e93cb82baa85bf759d93b7a7d4dde0..cd126f27f0a96372c78ed527fb429ff095a1b3a0 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/LBMKernel.cu
@@ -17,6 +17,8 @@
#include "GPU/GPU_Kernels.cuh"
#include "Parameter/Parameter.h"
+#include "Parameter/ParameterRotatingGrid.h"
+
//////////////////////////////////////////////////////////////////////////
void KernelCas27(
unsigned int grid_nx,
@@ -4052,6 +4054,70 @@ template<bool hasTurbulentViscosity> void ScaleCF_compressible(LBMSimulationPara
template void ScaleCF_compressible<true>(LBMSimulationParameter * parameterDeviceC, LBMSimulationParameter* parameterDeviceF, ICells * coarseToFine, ICellNeigh& neighborCoarseToFine, CUstream_st *stream);
template void ScaleCF_compressible<false>(LBMSimulationParameter * parameterDeviceC, LBMSimulationParameter* parameterDeviceF, ICells * coarseToFine, ICellNeigh& neighborCoarseToFine, CUstream_st *stream);
+void InterpolateStaticToRotating(
+ LBMSimulationParameter *parameterDeviceS,
+ LBMSimulationParameter *parameterDeviceR,
+ ParameterRotatingGridSimulation *paraRotDevice,
+ ICells *baseToNested,
+ ICellNeigh &neighborBaseToNested)
+{
+ dim3 grid = vf::cuda::getCudaGrid(parameterDeviceS->numberofthreads, baseToNested->numberOfCells);
+ dim3 threads(parameterDeviceS->numberofthreads, 1, 1);
+
+ interpolateStaticToRotating<<<grid, threads, 0, CU_STREAM_LEGACY>>>(
+ baseToNested->numberOfCells,
+ baseToNested->coarseCellIndices,
+ baseToNested->fineCellIndices,
+ paraRotDevice->nestedCoordinatesX,
+ paraRotDevice->nestedCoordinatesY,
+ paraRotDevice->nestedCoordinatesZ,
+ parameterDeviceS->coordinateX,
+ parameterDeviceS->coordinateY,
+ parameterDeviceS->coordinateZ,
+ parameterDeviceS->neighborX,
+ parameterDeviceS->neighborY,
+ parameterDeviceS->neighborZ,
+ parameterDeviceS->neighborInverse,
+ paraRotDevice->centerPoint[0],
+ paraRotDevice->centerPoint[1],
+ paraRotDevice->centerPoint[2],
+ paraRotDevice->gridAngle[0],
+ paraRotDevice->gridAngle[1],
+ paraRotDevice->gridAngle[2],
+ paraRotDevice->angularVelocity[0],
+ paraRotDevice->angularVelocity[1],
+ paraRotDevice->angularVelocity[2],
+ parameterDeviceS->gridSpacing
+ );
+
+ getLastCudaError("interpolateStaticToRotating execution failed");
+}
+
+void UpdateGlobalCoordinates(
+ LBMSimulationParameter *parameterDeviceR,
+ ParameterRotatingGridSimulation *paraRotDevice
+)
+{
+ dim3 grid = vf::cuda::getCudaGrid(parameterDeviceR->numberofthreads, parameterDeviceR->numberOfNodes);
+ dim3 threads(parameterDeviceR->numberofthreads, 1, 1);
+
+ updateGlobalCoordinates<<<grid, threads, 0, CU_STREAM_LEGACY>>>(
+ parameterDeviceR->numberOfNodes,
+ parameterDeviceR->coordinateX,
+ parameterDeviceR->coordinateY,
+ parameterDeviceR->coordinateZ,
+ paraRotDevice->nestedCoordinatesX,
+ paraRotDevice->nestedCoordinatesY,
+ paraRotDevice->nestedCoordinatesZ,
+ paraRotDevice->centerPoint[0],
+ paraRotDevice->centerPoint[1],
+ paraRotDevice->centerPoint[2],
+ paraRotDevice->gridAngle[0],
+ paraRotDevice->gridAngle[1],
+ paraRotDevice->gridAngle[2]
+ );
+ getLastCudaError("updateGlobalCoordinates execution failed");
+}
//////////////////////////////////////////////////////////////////////////
void ScaleCF_RhoSq_3rdMom_comp_27(
real* DC,
diff --git a/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/CoordinateTransformation.h b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/CoordinateTransformation.h
index 3c231bd2258bd09efa43e0a2300d176a7aa2916d..37967c012da12d2d4fdaf0c4a64ca6de8a3a7803 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/CoordinateTransformation.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/CoordinateTransformation.h
@@ -4,7 +4,7 @@
#include <basics/DataTypes.h>
#include <math.h>
-__inline__ void transformRotatingToGlobal(real &globalX, real &globalY, real &globalZ, real localX,
+__inline__ __device__ void transformRotatingToGlobal(real &globalX, real &globalY, real &globalZ, real localX,
real localY, real localZ, real centerCoordX, real centerCoordY,
real centerCoordZ, real angleX, real angleY, real angleZ)
{
diff --git a/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/GridTraversion.h b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/GridTraversion.h
new file mode 100644
index 0000000000000000000000000000000000000000..09a9fb5999bdb38f576792a2722e7a87b93173c8
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/GridTraversion.h
@@ -0,0 +1,30 @@
+#ifndef GRID_TRAVERSION
+#define GRID_TRAVERSION
+
+#include <basics/DataTypes.h>
+#include <math.h>
+
+__inline__ __host__ __device__ uint traverseSourceCell(real coordDestinationX, real coordDestinationY, real coordDestinationZ,
+ uint indexOfInverseNeighborOfSourceCell, real coordInvNeighborCellX, real coordInvNeighborCellY,
+ real coordInvNeighborCellZ, const uint *neighborX, const uint *neighborY, const uint *neighborZ, real dx)
+{
+ uint xTraverse = (uint)floor((coordDestinationX - coordInvNeighborCellX) / dx);
+ uint yTraverse = (uint)floor((coordDestinationY - coordInvNeighborCellY) / dx);
+ uint zTraverse = (uint)floor((coordDestinationZ - coordInvNeighborCellZ) / dx);
+
+ printf("s %.4f, d %.4f, delta %.4f, dx %.4f, xTraverse %d\n", coordInvNeighborCellX, coordDestinationX, coordDestinationX - coordInvNeighborCellX, dx, xTraverse);
+
+ uint newIndexOfSourceCell = indexOfInverseNeighborOfSourceCell;
+ for (uint ix = 1; ix <= xTraverse; ix++) {
+ newIndexOfSourceCell = neighborX[newIndexOfSourceCell];
+ }
+ for (uint iy = 1; iy <= yTraverse; iy++) {
+ newIndexOfSourceCell = neighborY[newIndexOfSourceCell];
+ }
+ for (uint iz = 1; iz <= zTraverse; iz++) {
+ newIndexOfSourceCell = neighborZ[newIndexOfSourceCell];
+ }
+ return newIndexOfSourceCell;
+}
+
+#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 b9d469d0e39d15413ff397faaadd400f3aa6f480..ac5cd5c6d2079a4c19adb082cd182ead10a526ad 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -10,6 +10,7 @@
#include "Parameter/Parameter.h"
#include "Parameter/CudaStreamManager.h"
#include "Parameter/EdgeNodeFinder.h"
+#include "Parameter/ParameterRotatingGrid.h"
#include "GPU/GPU_Interface.h"
#include "GPU/KineticEnergyAnalyzer.h"
#include "GPU/EnstrophyAnalyzer.h"
@@ -371,6 +372,7 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
dataWriter->writeInit(para, cudaMemoryManager);
if (para->getCalcParticles())
copyAndPrintParticles(para.get(), cudaMemoryManager.get(), 0, true);
+
VF_LOG_INFO("... done.");
//////////////////////////////////////////////////////////////////////////
@@ -381,8 +383,8 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
// NeighborDebugWriter::writeNeighborLinkLinesDebug(para.get());
- // InterfaceDebugWriter::writeInterfaceLinesDebugCF(para.get());
- // InterfaceDebugWriter::writeInterfaceLinesDebugFC(para.get());
+ InterfaceDebugWriter::writeInterfaceLinesDebugCF(para.get(), 0);
+ InterfaceDebugWriter::writeInterfaceLinesDebugFC(para.get(), 0);
// writers for version with communication hiding
// if(para->getNumprocs() > 1 && para->getUseStreams()){
@@ -965,6 +967,18 @@ void Simulation::readAndWriteFiles(uint timestep)
//writeAllTiDatafToFile(para.get(), t);
}
////////////////////////////////////////////////////////////////////////
+ // rotating grid
+
+ UpdateGlobalCoordinates(para->getParD(1).get(), para->getRotatingGridParameter()->parameterRotDevice.get());
+ cudaMemoryManager->cudaCopyCoordDeviceToHost(1);
+ // cudaMemoryManager->cudaCopyCoordRotationDeviceToHost(1);
+ cudaMemoryManager->cudaCopyInterfaceCFDeviceToHost(0);
+ cudaMemoryManager->cudaCopyInterfaceFCDeviceToHost(0);
+ InterfaceDebugWriter::writeInterfaceLinesDebugCF(para.get(), timestep);
+ InterfaceDebugWriter::writeInterfaceLinesDebugFC(para.get(), timestep);
+
+ ////////////////////////////////////////////////////////////////////////
+
dataWriter->writeTimestep(para, timestep);
////////////////////////////////////////////////////////////////////////
if (para->getCalcTurbulenceIntensity()) {
diff --git a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.cpp b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.cpp
index 5e96db99e17a0908ed9549af3c89ee48cf9db8d3..ffebe4053b0d05c35498e33917028cf65a0b15c0 100644
--- a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.cpp
+++ b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.cpp
@@ -39,18 +39,18 @@ void writeGridInterfaceLines(Parameter *para, int level, const uint *coarse, con
WbWriterVtkXmlBinary::getInstance()->writeLines(name, nodes, cells);
}
-void writeInterfaceLinesDebugCF(Parameter *para)
+void writeInterfaceLinesDebugCF(Parameter *para, uint timeStep)
{
for (int level = 0; level < para->getMaxLevel(); level++) {
- const std::string fileName = para->getFName() + "_" + StringUtil::toString<int>(level) + "_OffDebugCF.vtk";
+ const std::string fileName = para->getFName() + "_" + StringUtil::toString<int>(level) + "_OffDebugCF_" + StringUtil::toString<int>(timeStep) + ".vtk";
writeGridInterfaceLines(para, level, para->getParH(level)->coarseToFine.coarseCellIndices, para->getParH(level)->coarseToFine.fineCellIndices, para->getParH(level)->coarseToFine.numberOfCells, fileName);
}
}
-void writeInterfaceLinesDebugFC(Parameter *para)
+void writeInterfaceLinesDebugFC(Parameter *para, uint timeStep)
{
for (int level = 0; level < para->getMaxLevel(); level++) {
- const std::string fileName = para->getFName() + "_" + StringUtil::toString<int>(level) + "_OffDebugFC.vtk";
+ const std::string fileName = para->getFName() + "_" + StringUtil::toString<int>(level) + "_OffDebugFC_" + StringUtil::toString<int>(timeStep) + ".vtk";
writeGridInterfaceLines(para, level, para->getParH(level)->fineToCoarse.coarseCellIndices, para->getParH(level)->fineToCoarse.fineCellIndices, para->getParH(level)->fineToCoarse.numberOfCells, fileName);
}
}
diff --git a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.h b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.h
index 6a13f9bb48f5db7aced0009f58c87d773422f9c1..db1d0c44f09e2222a8695ba6fcf8755f0b3cde1b 100644
--- a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.h
+++ b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.h
@@ -1,12 +1,14 @@
#ifndef INTERFACEDEBUG_HPP
#define INTERFACEDEBUG_HPP
+#include <basics/DataTypes.h>
+
class Parameter;
namespace InterfaceDebugWriter
{
-void writeInterfaceLinesDebugCF(Parameter *para);
-void writeInterfaceLinesDebugFC(Parameter *para);
+void writeInterfaceLinesDebugCF(Parameter *para, uint timeStep=0);
+void writeInterfaceLinesDebugFC(Parameter *para, uint timeStep=0);
void writeInterfaceLinesDebugCFCneighbor(Parameter *para);
void writeInterfaceLinesDebugCFFneighbor(Parameter *para);
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
index 3b4a3bc25fe059c0a8149d3de97141a2c708490d..5c9f9b3411272b34e4544287eba416a38689067a 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp
@@ -518,7 +518,7 @@ void Parameter::initLBMSimulationParameter()
parH[i]->Lx = ((real)1.0 * parH[i]->gridNX - (real)1.0) / (pow((real)2.0, i));
parH[i]->Ly = ((real)1.0 * parH[i]->gridNY - (real)1.0) / (pow((real)2.0, i));
parH[i]->Lz = ((real)1.0 * parH[i]->gridNZ - (real)1.0) / (pow((real)2.0, i));
- parH[i]->dx = (real)1.0 / pow((real)2.0, i);
+ parH[i]->scalingFactorOfGridSpacing = (real)1.0 / pow((real)2.0, i);
parH[i]->XdistKn = getDistX().at(i);
parH[i]->YdistKn = getDistY().at(i);
parH[i]->ZdistKn = getDistZ().at(i);
@@ -539,9 +539,9 @@ void Parameter::initLBMSimulationParameter()
////////////////////////////////////////////////////////////////////////////
} else {
// Geller
- parH[i]->distX = ((real)getDistX().at(i) + (real)0.25) * parH[i - 1]->dx;
- parH[i]->distY = ((real)getDistY().at(i) + (real)0.25) * parH[i - 1]->dx;
- parH[i]->distZ = ((real)getDistZ().at(i) + (real)0.25) * parH[i - 1]->dx;
+ parH[i]->distX = ((real)getDistX().at(i) + (real)0.25) * parH[i - 1]->scalingFactorOfGridSpacing;
+ parH[i]->distY = ((real)getDistY().at(i) + (real)0.25) * parH[i - 1]->scalingFactorOfGridSpacing;
+ parH[i]->distZ = ((real)getDistZ().at(i) + (real)0.25) * parH[i - 1]->scalingFactorOfGridSpacing;
// parH[i]->distX = ((real)getDistX().at(i) + 0.25f) * parH[i-1]->dx + parH[i-1]->distX;
// parH[i]->distY = ((real)getDistY().at(i) + 0.25f) * parH[i-1]->dx + parH[i-1]->distY;
// parH[i]->distZ = ((real)getDistZ().at(i) + 0.25f) * parH[i-1]->dx + parH[i-1]->distZ;
@@ -586,7 +586,7 @@ void Parameter::initLBMSimulationParameter()
parD[i]->Lx = parH[i]->Lx;
parD[i]->Ly = parH[i]->Ly;
parD[i]->Lz = parH[i]->Lz;
- parD[i]->dx = parH[i]->dx;
+ parD[i]->scalingFactorOfGridSpacing = parH[i]->scalingFactorOfGridSpacing;
parD[i]->XdistKn = parH[i]->XdistKn;
parD[i]->YdistKn = parH[i]->YdistKn;
parD[i]->ZdistKn = parH[i]->ZdistKn;
@@ -1661,9 +1661,12 @@ void Parameter::setADKernel(std::string adKernel)
void Parameter::setRotatingGridParameter(std::shared_ptr<ParameterRotatingGrid> parameterRotatingGrid)
{
this->parameterRotatingGrid = std::move(parameterRotatingGrid);
- this->parameterRotatingGrid->initializeNestedCoordinates(
- { this->parH[1]->coordinateX, this->parH[1]->coordinateY, this->parH[1]->coordinateZ },
- this->parH[1]->numberOfNodes);
+}
+
+ void Parameter::fillCoordinateVectorsForRotatingGrid(uint level)
+{
+ this->parameterRotatingGrid->fillNestedCoordinateVectorsOnHost(
+ { this->parH[level]->coordinateX, this->parH[level]->coordinateY, this->parH[level]->coordinateZ });
}
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -2669,27 +2672,27 @@ real Parameter::TrafoXtoMGsWorld(int CoordX, int level)
{
real temp = 0;
for (int i = 0; i <= level; i++) {
- temp += (parH[i]->XdistKn + 0.25f) * 2.f * parH[i]->dx;
+ temp += (parH[i]->XdistKn + 0.25f) * 2.f * parH[i]->scalingFactorOfGridSpacing;
}
- temp += (real)((CoordX)*parH[level]->dx);
+ temp += (real)((CoordX)*parH[level]->scalingFactorOfGridSpacing);
return temp;
}
real Parameter::TrafoYtoMGsWorld(int CoordY, int level)
{
real temp = 0;
for (int i = 0; i <= level; i++) {
- temp += (parH[i]->YdistKn + 0.25f) * 2.f * parH[i]->dx;
+ temp += (parH[i]->YdistKn + 0.25f) * 2.f * parH[i]->scalingFactorOfGridSpacing;
}
- temp += (real)((CoordY)*parH[level]->dx);
+ temp += (real)((CoordY)*parH[level]->scalingFactorOfGridSpacing);
return temp;
}
real Parameter::TrafoZtoMGsWorld(int CoordZ, int level)
{
real temp = 0;
for (int i = 0; i <= level; i++) {
- temp += (parH[i]->ZdistKn + 0.25f) * 2.f * parH[i]->dx;
+ temp += (parH[i]->ZdistKn + 0.25f) * 2.f * parH[i]->scalingFactorOfGridSpacing;
}
- temp += (real)((CoordZ)*parH[level]->dx);
+ temp += (real)((CoordZ)*parH[level]->scalingFactorOfGridSpacing);
return temp;
}
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 9aba51d982f730abf86662f5c02ba4307ff1dd1b..82debe0aa1e9f11b3466ff969a2d7826c78c7d5d 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -441,13 +441,13 @@ struct LBMSimulationParameter {
// print///////////////////
unsigned int startz, endz;
- real Lx, Ly, Lz, dx;
+ real Lx, Ly, Lz, scalingFactorOfGridSpacing;
real distX, distY, distZ;
// testRoundoffError
Distributions27 kDistTestRE;
-
+ real gridSpacing;
//////////////////////////////////////////////////////////////////////////
};
@@ -686,6 +686,7 @@ public:
void setADKernel(std::string adKernel);
// rotating grid
void setRotatingGridParameter(std::shared_ptr<ParameterRotatingGrid> parameterRotatingGrid);
+ void fillCoordinateVectorsForRotatingGrid(uint level);
// adder
void addActuator(SPtr<PreCollisionInteractor> actuator);
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.cpp b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.cpp
index dca388592516994ac2b5e7c80034b4bd30167cd5..5031aeb145d4338ee4de8b195dd892de2f36dfa0 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.cpp
+++ b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.cpp
@@ -1,33 +1,40 @@
#include "ParameterRotatingGrid.h"
-#include "LBM/GPUHelperFunctions/CoordinateTransformation.h"
+#include "Logger.h"
+#include <iostream>
-ParameterRotatingGrid::ParameterRotatingGrid(const std::array<real, 3> ¢erPoint, const Axis &rotationalAxis)
- : centerPoint(centerPoint), rotationalAxis(rotationalAxis)
+ParameterRotatingGrid::ParameterRotatingGrid(const std::array<real, 3> ¢erPoint, const Axis &rotationalAxis,
+ uint sizeOfNestedCoordinates)
+ : rotationalAxis(rotationalAxis)
{
+ parameterRotHost->centerPoint = centerPoint;
+ parameterRotHost->sizeOfNestedCoordinates = sizeOfNestedCoordinates;
+ parameterRotHost->memorySizeOfNestedCoordinates = sizeOfNestedCoordinates * sizeof(real);
+
+ parameterRotDevice->centerPoint = centerPoint;
+ parameterRotDevice->sizeOfNestedCoordinates = sizeOfNestedCoordinates;
+ parameterRotDevice->memorySizeOfNestedCoordinates = sizeOfNestedCoordinates * sizeof(real);
}
-void ParameterRotatingGrid::initializeNestedCoordinates(const std::array<real *, 3> &globalCoordinates, uint numberOfNodes)
+void ParameterRotatingGrid::fillNestedCoordinateVectorsOnHost(const std::array<real *, 3> &globalCoordinates)
{
- this->nestedCoordinatesX.resize(numberOfNodes);
- this->nestedCoordinatesY.resize(numberOfNodes);
- this->nestedCoordinatesZ.resize(numberOfNodes);
+ if (parameterRotHost->nestedCoordinatesX == nullptr || parameterRotHost->nestedCoordinatesY == nullptr ||
+ parameterRotHost->nestedCoordinatesZ == nullptr)
+ throw std::runtime_error("Allocate host pointers of nestedCoordinates before filing them!");
-// #pragma omp parallel for
- for (uint index = 0; index < numberOfNodes; index++) {
- this->nestedCoordinatesX[index] = globalCoordinates[0][index] - centerPoint[0];
- this->nestedCoordinatesY[index] = globalCoordinates[1][index] - centerPoint[1];
- this->nestedCoordinatesZ[index] = globalCoordinates[2][index] - centerPoint[2];
+ // #pragma omp parallel for
+ for (uint index = 0; index < parameterRotHost->sizeOfNestedCoordinates; index++) {
+ this->parameterRotHost->nestedCoordinatesX[index] = globalCoordinates[0][index] - parameterRotHost->centerPoint[0];
+ this->parameterRotHost->nestedCoordinatesY[index] = globalCoordinates[1][index] - parameterRotHost->centerPoint[1];
+ this->parameterRotHost->nestedCoordinatesZ[index] = globalCoordinates[2][index] - parameterRotHost->centerPoint[2];
}
}
-void ParameterRotatingGrid::transformNestedToBase(const std::array<real *, 3> &globalCoordinates)
-{
-// #pragma omp parallel for
- for (uint index = 0; index < nestedCoordinatesX.size(); index++) {
- transformRotatingToGlobal(
- globalCoordinates[0][index], globalCoordinates[1][index], globalCoordinates[2][index], nestedCoordinatesX[index],
- nestedCoordinatesY[index], nestedCoordinatesZ[index], centerPoint[0], centerPoint[1], centerPoint[2],
- gridAngle * unitVectors.at(this->rotationalAxis)[0], gridAngle * unitVectors.at(this->rotationalAxis)[1],
- gridAngle * unitVectors.at(this->rotationalAxis)[2]);
- }
-}
+// void ParameterRotatingGrid::transformNestedToBase(const std::array<real *, 3> &globalCoordinates)
+// {
+// // #pragma omp parallel for
+// for (uint index = 0; index < nestedCoordinatesX.size(); index++) {
+// transformRotatingToGlobal(globalCoordinates[0][index], globalCoordinates[1][index], globalCoordinates[2][index],
+// nestedCoordinatesX[index], nestedCoordinatesY[index], nestedCoordinatesZ[index],
+// centerPoint[0], centerPoint[1], centerPoint[2], gridAngle[0], gridAngle[1], gridAngle[2]);
+// }
+// }
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.h b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.h
index 573c975e7df7acdef6e949aa96a3b13e8fa5be22..ac78941f87564936acca189f8d3f3ae81716ca0c 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGrid.h
@@ -1,28 +1,32 @@
#ifndef PARAMETER_ROTATING_GRID
#define PARAMETER_ROTATING_GRID
+#include "PointerDefinitions.h"
#include <array>
#include <vector>
#include <basics/DataTypes.h>
#include <basics/geometry3d/Axis.h>
+struct ParameterRotatingGridSimulation {
+ std::array<real, 3> centerPoint;
+ real *nestedCoordinatesX = nullptr; // in local coordinate system of rotating grid
+ real *nestedCoordinatesY = nullptr;
+ real *nestedCoordinatesZ = nullptr;
+ uint sizeOfNestedCoordinates;
+ uint memorySizeOfNestedCoordinates;
+ std::array<real, 3> gridAngle = { 0.0, 0.0, 0.0 };
+ std::array<real, 3> angularVelocity = { 0.02, 0.0, 0.0 };
+};
+
struct ParameterRotatingGrid {
public:
- ParameterRotatingGrid(const std::array<real, 3> ¢erPoint, const Axis &rotationalAxis);
- void initializeNestedCoordinates(const std::array<real *, 3> &globalCoordinates, uint numberOfNodes);
- void transformNestedToBase(const std::array<real *, 3> &globalCoordinates);
+ ParameterRotatingGrid(const std::array<real, 3> ¢erPoint, const Axis &rotationalAxis, uint sizeOfNestedCoordinates);
+ void fillNestedCoordinateVectorsOnHost(const std::array<real *, 3> &globalCoordinates);
-public:
- const std::array<real, 3> centerPoint;
const Axis rotationalAxis;
-
- std::vector<real> nestedCoordinatesX;
- std::vector<real> nestedCoordinatesY;
- std::vector<real> nestedCoordinatesZ;
-
- real gridAngle = 1.0;
- std::array<real, 3> angularVelocity;
+ SPtr<ParameterRotatingGridSimulation> parameterRotHost = std::make_shared<ParameterRotatingGridSimulation>();
+ SPtr<ParameterRotatingGridSimulation> parameterRotDevice = std::make_shared<ParameterRotatingGridSimulation>();
};
#endif
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGridTest.cpp b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGridTest.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..bdd77dd6386a46d1477bc80f3211d75c2c6111af
--- /dev/null
+++ b/src/gpu/VirtualFluids_GPU/Parameter/ParameterRotatingGridTest.cpp
@@ -0,0 +1,6 @@
+#include <gmock/gmock.h>
+
+// TEST(ParameterRotatingGrid, arrayForCoordinatesHostNotAllocated_fillWithCoordinates_Throws)
+// {
+// CudaMemoryManager cuda
+// }
\ No newline at end of file