From b720b4257d79e06a14ad7a06f4c85e46547ed1f1 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-braunschweig.de>
Date: Thu, 28 Sep 2023 10:49:42 +0000
Subject: [PATCH] Fix problems after merge
---
apps/gpu/RotatingGrid/RotatingGrid.cpp | 2 +-
src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu | 83 ++++++++-----------
.../interpolateRotatingToStatic.cu | 4 +-
.../interpolateStaticToRotating.cu | 4 +-
.../LBM/GPUHelperFunctions/KernelUtilities.h | 1 +
5 files changed, 40 insertions(+), 54 deletions(-)
diff --git a/apps/gpu/RotatingGrid/RotatingGrid.cpp b/apps/gpu/RotatingGrid/RotatingGrid.cpp
index 76061a25f..c3db2a289 100644
--- a/apps/gpu/RotatingGrid/RotatingGrid.cpp
+++ b/apps/gpu/RotatingGrid/RotatingGrid.cpp
@@ -141,7 +141,7 @@ int main()
para->setTimestepEnd(timeStepEnd);
para->setTimestepStartOut(timeStepStartOutput);
- para->configureMainKernel(vf::CollisionKernel::Compressible::K15CompressibleNavierStokes);
+ para->configureMainKernel(vf::collisionKernel::compressible::K15CompressibleNavierStokes);
//////////////////////////////////////////////////////////////////////////
// set boundary conditions
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu
index 437054955..201464979 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/CalcMac27.cu
@@ -272,6 +272,33 @@ __global__ void LBCalcMacSP27(
+__inline__ __device__ void calcMacCompSP27ForNode(real* vxD, real* vyD, real* vzD, real* rhoD, real* pressD,
+ unsigned int* geoD, unsigned int* neighborX, unsigned int* neighborY,
+ unsigned int* neighborZ, unsigned long long numberOfLBnodes,
+ real* distributions, bool isEvenTimestep, unsigned int nodeIndex)
+{
+ pressD[nodeIndex] = c0o1;
+ rhoD[nodeIndex] = c0o1;
+ vxD[nodeIndex] = c0o1;
+ vyD[nodeIndex] = c0o1;
+ vzD[nodeIndex] = c0o1;
+
+ if (!isValidFluidNode(geoD[nodeIndex]))
+ return;
+
+ Distributions27 dist;
+ vf::gpu::getPointersToDistributions(dist, distributions, numberOfLBnodes, isEvenTimestep);
+ vf::gpu::ListIndices listIndices(nodeIndex, neighborX, neighborY, neighborZ);
+
+ real distribution[27];
+ vf::gpu::read(distribution, dist, listIndices);
+
+ rhoD[nodeIndex] = vf::lbm::getDensity(distribution);
+ vxD[nodeIndex] = vf::lbm::getCompressibleVelocityX1(distribution, rhoD[nodeIndex]);
+ vyD[nodeIndex] = vf::lbm::getCompressibleVelocityX2(distribution, rhoD[nodeIndex]);
+ vzD[nodeIndex] = vf::lbm::getCompressibleVelocityX3(distribution, rhoD[nodeIndex]);
+ pressD[nodeIndex] = vf::lbm::getPressure(distribution, rhoD[nodeIndex], vxD[nodeIndex], vyD[nodeIndex], vzD[nodeIndex]);
+}
////////////////////////////////////////////////////////////////////////////////
__global__ void LBCalcMacCompSP27(
@@ -288,35 +315,14 @@ __global__ void LBCalcMacCompSP27(
real *distributions,
bool isEvenTimestep)
{
- ////////////////////////////////////////////////////////////////////////////////
- //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
- //!
const unsigned nodeIndex = getNodeIndex();
if(nodeIndex >= numberOfLBnodes)
return;
- pressD[nodeIndex] = c0o1;
- rhoD[nodeIndex] = c0o1;
- vxD[nodeIndex] = c0o1;
- vyD[nodeIndex] = c0o1;
- vzD[nodeIndex] = c0o1;
-
- if (!isValidFluidNode(geoD[nodeIndex]))
- return;
-
- Distributions27 dist;
- vf::gpu::getPointersToDistributions(dist, distributions, numberOfLBnodes, isEvenTimestep);
- vf::gpu::ListIndices listIndices(nodeIndex, neighborX, neighborY, neighborZ);
-
- real distribution[27];
- vf::gpu::read(distribution, dist, listIndices);
-
- rhoD[nodeIndex] = vf::lbm::getDensity(distribution);
- vxD[nodeIndex] = vf::lbm::getCompressibleVelocityX1(distribution, rhoD[nodeIndex]);
- vyD[nodeIndex] = vf::lbm::getCompressibleVelocityX2(distribution, rhoD[nodeIndex]);
- vzD[nodeIndex] = vf::lbm::getCompressibleVelocityX3(distribution, rhoD[nodeIndex]);
- pressD[nodeIndex] = vf::lbm::getPressure(distribution, rhoD[nodeIndex], vxD[nodeIndex], vyD[nodeIndex], vzD[nodeIndex]);
+ calcMacCompSP27ForNode(vxD, vyD, vzD, rhoD, pressD, geoD,
+ neighborX, neighborY, neighborZ,
+ numberOfLBnodes, distributions, isEvenTimestep, nodeIndex);
}
@@ -339,37 +345,16 @@ __global__ void LBCalcMacCompSP27RotatingToStatic(
real angleZ
)
{
- ////////////////////////////////////////////////////////////////////////////////
- //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim.
- //!
const unsigned nodeIndex = getNodeIndex();
if(nodeIndex >= numberOfLBnodes)
return;
- pressD[nodeIndex] = c0o1;
- rhoD[nodeIndex] = c0o1;
- vxD[nodeIndex] = c0o1;
- vyD[nodeIndex] = c0o1;
- vzD[nodeIndex] = c0o1;
-
- if (!isValidFluidNode(geoD[nodeIndex]))
- return;
-
- DistributionWrapper distr_wrapper(distributions, numberOfLBnodes, isEvenTimestep, nodeIndex, neighborX, neighborY, neighborZ);
- const auto &distribution = distr_wrapper.distribution;
-
- rhoD[nodeIndex] = vf::lbm::getDensity(distribution.f);
- real velocityRotatingX = vf::lbm::getCompressibleVelocityX1(distribution.f, rhoD[nodeIndex]);
- real velocityRotatingY = vf::lbm::getCompressibleVelocityX2(distribution.f, rhoD[nodeIndex]);
- real velocityRotatingZ = vf::lbm::getCompressibleVelocityX3(distribution.f, rhoD[nodeIndex]);
- pressD[nodeIndex] = vf::lbm::getPressure(distribution.f, rhoD[nodeIndex], vxD[nodeIndex], vyD[nodeIndex], vzD[nodeIndex]);
-
- rotateDataFromGlobalToRotating(velocityRotatingX, velocityRotatingY, velocityRotatingZ, angleX, angleY, angleZ);
+ calcMacCompSP27ForNode(vxD, vyD, vzD, rhoD, pressD, geoD,
+ neighborX, neighborY, neighborZ,
+ numberOfLBnodes, distributions, isEvenTimestep, nodeIndex);
- vxD[nodeIndex] = velocityRotatingX;
- vyD[nodeIndex] = velocityRotatingY;
- vzD[nodeIndex] = velocityRotatingZ;
+ rotateDataFromGlobalToRotating(vxD[nodeIndex], vyD[nodeIndex], vzD[nodeIndex], angleX, angleY, angleZ);
}
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateRotatingToStatic.cu b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateRotatingToStatic.cu
index 236d8c8ab..b2d86b05d 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateRotatingToStatic.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateRotatingToStatic.cu
@@ -40,7 +40,7 @@
#include "GPU/GridScaling/interpolateRotatingToStaticInlines.h"
#include <lbm/MacroscopicQuantities.h>
-#include <lbm/refinement/Coefficients.h>
+#include <lbm/interpolation/InterpolationCoefficients.h>
#include <lbm/refinement/InterpolationCF.h>
@@ -289,7 +289,7 @@ __global__ void interpolateRotatingToStatic(
const real cellCoordDestinationY = (coordSourceY[sourceIndex] - rotatedCoordDestinationY) / dx + c1o2;
const real cellCoordDestinationZ = (coordSourceZ[sourceIndex] - rotatedCoordDestinationZ) / dx + c1o2;
momentsSet.addToVelocity(tangentialVelocitiesX, tangentialVelocitiesY, tangentialVelocitiesZ);
- vf::lbm::Coefficients coefficients;
+ vf::lbm::InterpolationCoefficients coefficients;
momentsSet.calculateCoefficients(coefficients, cellCoordDestinationX, cellCoordDestinationY, cellCoordDestinationZ);
////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticToRotating.cu b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticToRotating.cu
index 7f8af0357..6e80893b4 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticToRotating.cu
+++ b/src/gpu/VirtualFluids_GPU/GPU/GridScaling/interpolateStaticToRotating.cu
@@ -39,7 +39,7 @@
#include "LBM/GPUHelperFunctions/CoordinateTransformation.h"
#include "GPU/GridScaling/interpolateStaticToRotatingInlines.h"
-#include <lbm/refinement/Coefficients.h>
+#include <lbm/interpolation/InterpolationCoefficients.h>
#include <lbm/refinement/InterpolationCF.h>
using namespace vf::lbm;
@@ -117,7 +117,7 @@ __global__ void interpolateStaticToRotating(
const real cellCoordDestinationX = (coordSourceX[sourceIndex] - globalCoordDestinationX) / dx + c1o2;
const real cellCoordDestinationY = (coordSourceY[sourceIndex] - globalCoordDestinationY) / dx + c1o2;
const real cellCoordDestinationZ = (coordSourceZ[sourceIndex] - globalCoordDestinationZ) / dx + c1o2;
- vf::lbm::Coefficients coefficients;
+ vf::lbm::InterpolationCoefficients coefficients;
momentsSet.calculateCoefficients(coefficients, cellCoordDestinationX, cellCoordDestinationY, cellCoordDestinationZ);
////////////////////////////////////////////////////////////////////////////////////
diff --git a/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/KernelUtilities.h b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/KernelUtilities.h
index 31aaf6eb3..f67a2d752 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/KernelUtilities.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/GPUHelperFunctions/KernelUtilities.h
@@ -213,6 +213,7 @@ __inline__ __device__ real getInterpolatedDistributionForVeloWithPressureBC(cons
__inline__ __device__ unsigned int getNodeIndex()
{
+ // get node index from CUDA threadIdx, blockIdx, blockDim and gridDim.
const unsigned x = threadIdx.x;
const unsigned y = blockIdx.x;
const unsigned z = blockIdx.y;
--
GitLab