diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
index af1e123e513e56dd4d500fee1f98bbd05c5f6240..48fe29321f5f884a2e73118bab1ea352fcf5f6ec 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridProvider.cpp
@@ -80,8 +80,8 @@ void GridProvider::setInitalNodeValues(const int numberOfNodes) const
void GridProvider::setVelocitySize(int size) const
{
- para->getParH()->numberOfInflowBCnodes = size;
- para->getParD()->numberOfInflowBCnodes = size;
+ para->getParH()->numberOfVeloBCnodes = size;
+ para->getParD()->numberOfVeloBCnodes = size;
}
void GridProvider::allocAndCopyForcing()
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index 75458df1c59a3a032b634cd12cd24ff2fd456a80..a68e8543f8673329f51be9de97d108e6eba4233f 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -119,20 +119,20 @@ void GridGenerator::allocArrays_BoundaryValues()
std::cout << "size velocity: " << numberOfVelocityValues << std::endl;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
blocks = (numberOfVelocityValues / para->getParH()->numberofthreads) + 1;
- para->getParH()->inflowBC.kArray = blocks * para->getParH()->numberofthreads;
- para->getParD()->inflowBC.kArray = para->getParH()->inflowBC.kArray;
+ para->getParH()->veloBC.kArray = blocks * para->getParH()->numberofthreads;
+ para->getParD()->veloBC.kArray = para->getParH()->veloBC.kArray;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- para->getParH()->numberOfInflowBCnodes = numberOfVelocityValues;
- para->getParD()->numberOfInflowBCnodes = numberOfVelocityValues;
+ para->getParH()->numberOfVeloBCnodes = numberOfVelocityValues;
+ para->getParD()->numberOfVeloBCnodes = numberOfVelocityValues;
if (numberOfVelocityValues > 1) {
cudaMemoryManager->cudaAllocVeloBC();
builder->getVelocityValues(
- para->getParH()->inflowBC.Vx,
- para->getParH()->inflowBC.Vy,
- para->getParH()->inflowBC.Vz,
- para->getParH()->inflowBC.k, 0);
+ para->getParH()->veloBC.Vx,
+ para->getParH()->veloBC.Vy,
+ para->getParH()->veloBC.Vz,
+ para->getParH()->veloBC.k, 0);
cudaMemoryManager->cudaCopyVeloBC();
}
@@ -192,8 +192,8 @@ void GridGenerator::allocArrays_BoundaryQs()
std::cout << "size velocity: " << numberOfVelocityNodes << std::endl;
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// preprocessing
- real *QQ = para->getParH()->inflowBC.q27[0];
- unsigned int sizeQ = para->getParH()->numberOfInflowBCnodes;
+ real *QQ = para->getParH()->veloBC.q27[0];
+ unsigned int sizeQ = para->getParH()->numberOfVeloBCnodes;
QforBoundaryConditions Q;
Q.q27[dirE] = &QQ[dirE * sizeQ];
Q.q27[dirW] = &QQ[dirW * sizeQ];
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
index 451174fa70031aca85c806f72f79e378d667220d..ee0dd86267e49733b8e35b70c7b9a7106120cf5c 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaKernelManager.cpp
@@ -69,18 +69,18 @@ void CudaKernelManager::runLBMKernel(SPtr<Parameter> para)
void CudaKernelManager::runVelocityBCKernel(SPtr<Parameter> para)
{
- if (para->getParD()->numberOfInflowBCnodes > 0)
+ if (para->getParD()->numberOfVeloBCnodes > 0)
{
QVelDevicePlainBB27(
para->getParD()->numberofthreads,
- para->getParD()->inflowBC.Vx,
- para->getParD()->inflowBC.Vy,
- para->getParD()->inflowBC.Vz,
+ para->getParD()->veloBC.Vx,
+ para->getParD()->veloBC.Vy,
+ para->getParD()->veloBC.Vz,
para->getParD()->distributions.f[0],
- para->getParD()->inflowBC.k,
- para->getParD()->inflowBC.q27[0],
- para->getParD()->numberOfInflowBCnodes,
- para->getParD()->inflowBC.kArray,
+ para->getParD()->veloBC.k,
+ para->getParD()->veloBC.q27[0],
+ para->getParD()->numberOfVeloBCnodes,
+ para->getParD()->veloBC.kArray,
para->getParD()->neighborX,
para->getParD()->neighborY,
para->getParD()->neighborZ,
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 73c6d1acc10119b1d3de46c540a5f8fa785fd4f8..364843deb94fd6730940055f5cba10179f057265 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -192,22 +192,22 @@ void CudaMemoryManager::cudaFreeSlipBC()
//velocity boundary condition
void CudaMemoryManager::cudaAllocVeloBC()
{
- unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfInflowBCnodes;
- unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfInflowBCnodes;
+ unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfVeloBCnodes;
+ unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfVeloBCnodes;
//Host
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.k), mem_size_inflow_BC_INT ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vx), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vy), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->inflowBC.Vz), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.k), mem_size_inflow_BC_INT ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vx), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vy), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH()->veloBC.Vz), mem_size_inflow_BC_REAL ));
//Device
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.k), mem_size_inflow_BC_INT ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vx), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vy), mem_size_inflow_BC_REAL ));
- checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->inflowBC.Vz), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.q27[0]), parameter->getD3Qxx() * mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.k), mem_size_inflow_BC_INT ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vx), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vy), mem_size_inflow_BC_REAL ));
+ checkCudaErrors( cudaMalloc((void**) &(parameter->getParD()->veloBC.Vz), mem_size_inflow_BC_REAL ));
//////////////////////////////////////////////////////////////////////////
double tmp = (double)mem_size_inflow_BC_INT + 4. * (double)mem_size_inflow_BC_REAL + (double)parameter->getD3Qxx() * (double)mem_size_inflow_BC_REAL;
@@ -215,23 +215,23 @@ void CudaMemoryManager::cudaAllocVeloBC()
}
void CudaMemoryManager::cudaCopyVeloBC()
{
- unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfInflowBCnodes;
- unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfInflowBCnodes;
+ unsigned int mem_size_inflow_BC_INT = sizeof(int)*parameter->getParH()->numberOfVeloBCnodes;
+ unsigned int mem_size_inflow_BC_REAL = sizeof(real)*parameter->getParH()->numberOfVeloBCnodes;
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.q27[0], parameter->getParH()->inflowBC.q27[0], parameter->getD3Qxx() * mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.k, parameter->getParH()->inflowBC.k, mem_size_inflow_BC_INT , cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vx, parameter->getParH()->inflowBC.Vx, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vy, parameter->getParH()->inflowBC.Vy, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
- checkCudaErrors( cudaMemcpy(parameter->getParD()->inflowBC.Vz, parameter->getParH()->inflowBC.Vz, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.q27[0], parameter->getParH()->veloBC.q27[0], parameter->getD3Qxx() * mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.k, parameter->getParH()->veloBC.k, mem_size_inflow_BC_INT , cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vx, parameter->getParH()->veloBC.Vx, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vy, parameter->getParH()->veloBC.Vy, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
+ checkCudaErrors( cudaMemcpy(parameter->getParD()->veloBC.Vz, parameter->getParH()->veloBC.Vz, mem_size_inflow_BC_REAL, cudaMemcpyHostToDevice));
}
void CudaMemoryManager::cudaFreeVeloBC()
{
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.q27[0] ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.k ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vx ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vy ));
- checkCudaErrors( cudaFreeHost(parameter->getParH()->inflowBC.Vz ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.q27[0] ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.k ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vx ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vy ));
+ checkCudaErrors( cudaFreeHost(parameter->getParH()->veloBC.Vz ));
}
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index 3e75d3c4cc29f8c963dfe4d2edc78fec884385d3..95d3f15dff5d3a31240c965a36d243096ba1a2ef 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -144,7 +144,7 @@ void Simulation::run()
////////////////////////////////////////////////////////////////////////////////
// velocity boundary condition
- if (para->getParD()->numberOfInflowBCnodes > 1)
+ if (para->getParD()->numberOfVeloBCnodes > 1)
cudaKernelManager->runVelocityBCKernel(para);
////////////////////////////////////////////////////////////////////////////////
@@ -194,7 +194,7 @@ void Simulation::free()
cudaMemoryManager->cudaFreeSP();
cudaMemoryManager->cudaFreeNeighborWSB();
- if (para->getParH()->numberOfInflowBCnodes > 1)
+ if (para->getParH()->numberOfVeloBCnodes > 1)
this->cudaMemoryManager->cudaFreeVeloBC();
if (para->getParH()->numberOfSlipBCnodes > 1)
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index 24b472f5ac3230e2925389d66b7a840f813f5e31..912c1c28314e66210ce5a293832be99563158dfd 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -80,9 +80,9 @@ struct ParameterStruct{
uint numberOfSlipBCnodes;
//////////////////////////////////////////////////////////////////////////
//! \brief stores the velocity boundary condition data
- QforBoundaryConditions inflowBC;
+ QforBoundaryConditions veloBC;
//! \brief number of lattice nodes for the velocity boundary condition
- uint numberOfInflowBCnodes;
+ uint numberOfVeloBCnodes;
//////////////////////////////////////////////////////////////////////////
//! \brief sets the forcing uniform on every fluid node in all three space dimensions
real *forcing;