From 2bc8cc80817fd9cc430916bc7fc49aef025a29b9 Mon Sep 17 00:00:00 2001
From: HenrikAsmuth <henrik.asmuth@geo.uu.se>
Date: Tue, 13 Dec 2022 16:29:11 +0100
Subject: [PATCH] Rename CollisionTemplate:: and CudaStreamIndex::Border to
SubDomainBorder
---
src/gpu/GridGenerator/grid/GridImp.h | 2 +-
.../Calculation/CollisisionStrategy.cpp | 12 ++++----
.../Calculation/RefinementStrategy.cpp | 28 +++++++++----------
.../Communication/ExchangeData27.cpp | 2 +-
.../GridReaderGenerator/GridGenerator.cpp | 16 +++++------
.../GPU/CudaMemoryManager.cpp | 24 ++++++++--------
.../Compressible/CumulantK17/CumulantK17.cu | 2 +-
src/gpu/VirtualFluids_GPU/LBM/LB.h | 4 +--
src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp | 2 +-
.../Parameter/CudaStreamManager.h | 2 +-
.../VirtualFluids_GPU/Parameter/Parameter.h | 4 +--
11 files changed, 49 insertions(+), 49 deletions(-)
diff --git a/src/gpu/GridGenerator/grid/GridImp.h b/src/gpu/GridGenerator/grid/GridImp.h
index d330bc8b8..8283bf569 100644
--- a/src/gpu/GridGenerator/grid/GridImp.h
+++ b/src/gpu/GridGenerator/grid/GridImp.h
@@ -116,7 +116,7 @@ private:
int *sparseIndices;
std::vector<uint> fluidNodeIndices; // run on CollisionTemplate::Default
- std::vector<uint> fluidNodeIndicesBorder; // run on subdomain border nodes (CollisionTemplate::Border)
+ std::vector<uint> fluidNodeIndicesBorder; // run on subdomain border nodes (CollisionTemplate::SubDomainBorder)
std::vector<uint> fluidNodeIndicesMacroVars; // run on CollisionTemplate::MacroVars
std::vector<uint> fluidNodeIndicesApplyBodyForce; // run on CollisionTemplate::ApplyBodyForce
std::vector<uint> fluidNodeIndicesAllFeatures; // run on CollisionTemplate::AllFeatures
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
index 1c279ad52..49543f37d 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/CollisisionStrategy.cpp
@@ -72,16 +72,16 @@ void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Paramete
//! 1. run collision for nodes which are at the border of the gpus/processes, running with WriteMacroVars in case probes sample on these nodes
//!
updateGrid->collisionUsingIndices( level, t,
- para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Border],
- para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Border],
+ para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::SubDomainBorder],
+ para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::SubDomainBorder],
CollisionTemplate::WriteMacroVars,
- CudaStreamIndex::Border);
+ CudaStreamIndex::SubDomainBorder);
//! 2. prepare the exchange between gpus (collect the send nodes for communication in a buffer on the gpu) and trigger bulk kernel execution when finished
//!
- updateGrid->prepareExchangeMultiGPU(level, CudaStreamIndex::Border);
+ updateGrid->prepareExchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
if (para->getUseStreams())
- para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::Border);
+ para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::SubDomainBorder);
//! 3. launch the collision kernel for bulk nodes. This includes nodes with \param tag Default, WriteMacroVars, ApplyBodyForce,
//! or AllFeatures. All assigned tags are listed in \param allocatedBulkFluidNodeTags during initialization in Simulation::init
@@ -97,5 +97,5 @@ void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Paramete
CudaStreamIndex::Bulk);
}
//! 4. exchange information between GPUs
- updateGrid->exchangeMultiGPU(level, CudaStreamIndex::Border);
+ updateGrid->exchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
}
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp
index fe7da04d4..b8ca4e9c2 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp
@@ -42,27 +42,27 @@ void RefinementAndExchange_streams_exchangeInterface::operator()(UpdateGrid27 *u
//!
//! 1. Interpolation fine to coarse for nodes which are at the border of the gpus/processes
//!
- updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBorder, para->getParD(level)->offFC, CudaStreamIndex::Border);
+ updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBorder, para->getParD(level)->offFC, CudaStreamIndex::SubDomainBorder);
//! 2. prepare the exchange between gpus (collect the send nodes for communication in a buffer on the gpu) and trigger bulk kernel execution when finished
//!
- updateGrid->prepareExchangeMultiGPUAfterFtoC(level, CudaStreamIndex::Border);
+ updateGrid->prepareExchangeMultiGPUAfterFtoC(level, CudaStreamIndex::SubDomainBorder);
if (para->getUseStreams())
- para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::Border);
+ para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::SubDomainBorder);
//! 3. launch the bulk kernels for both interpolation processes (fine to coarse and coarse to fine)
//!
para->getStreamManager()->waitOnStartBulkKernelEvent(CudaStreamIndex::Bulk);
- updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBulk, para->getParD(level)->offFCBulk, CudaStreamIndex::Border);
- updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBulk, para->getParD(level)->offCFBulk, CudaStreamIndex::Border);
+ updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBulk, para->getParD(level)->offFCBulk, CudaStreamIndex::SubDomainBorder);
+ updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBulk, para->getParD(level)->offCFBulk, CudaStreamIndex::SubDomainBorder);
//! 4. exchange information between GPUs (only nodes which are part of the interpolation)
//!
- updateGrid->exchangeMultiGPUAfterFtoC(level, CudaStreamIndex::Border);
+ updateGrid->exchangeMultiGPUAfterFtoC(level, CudaStreamIndex::SubDomainBorder);
// 5. interpolation fine to coarse for nodes which are at the border of the gpus/processes
//!
- updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBorder, para->getParD(level)->offCF, CudaStreamIndex::Border);
+ updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBorder, para->getParD(level)->offCF, CudaStreamIndex::SubDomainBorder);
cudaDeviceSynchronize();
}
@@ -73,27 +73,27 @@ void RefinementAndExchange_streams_exchangeAllNodes::operator()(UpdateGrid27 *up
//!
//! 1. interpolation fine to coarse for nodes which are at the border of the gpus/processes
//!
- updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBorder, para->getParD(level)->offFC, CudaStreamIndex::Border);
+ updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBorder, para->getParD(level)->offFC, CudaStreamIndex::SubDomainBorder);
//! 2. prepare the exchange between gpus (collect the send nodes for communication in a buffer on the gpu) and trigger bulk kernel execution when finished
//!
- updateGrid->prepareExchangeMultiGPU(level, CudaStreamIndex::Border);
+ updateGrid->prepareExchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
if (para->getUseStreams())
- para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::Border);
+ para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::SubDomainBorder);
//! 3. launch the bulk kernels for both interpolation processes (fine to coarse and coarse to fine)
//!
para->getStreamManager()->waitOnStartBulkKernelEvent(CudaStreamIndex::Bulk);
- updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBulk, para->getParD(level)->offFCBulk, CudaStreamIndex::Border);
- updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBulk, para->getParD(level)->offCFBulk, CudaStreamIndex::Border);
+ updateGrid->fineToCoarse(level, ¶->getParD(level)->intFCBulk, para->getParD(level)->offFCBulk, CudaStreamIndex::SubDomainBorder);
+ updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBulk, para->getParD(level)->offCFBulk, CudaStreamIndex::SubDomainBorder);
//! 4. exchange information between GPUs (all nodes)
//!
- updateGrid->exchangeMultiGPU(level, CudaStreamIndex::Border);
+ updateGrid->exchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
// 5. interpolation fine to coarse for nodes which are at the border of the gpus/processes
//!
- updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBorder, para->getParD(level)->offCF, CudaStreamIndex::Border);
+ updateGrid->coarseToFine(level, ¶->getParD(level)->intCFBorder, para->getParD(level)->offCF, CudaStreamIndex::SubDomainBorder);
cudaDeviceSynchronize();
}
diff --git a/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp b/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
index 44202b716..00a7b4566 100644
--- a/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Communication/ExchangeData27.cpp
@@ -240,7 +240,7 @@ void exchangeCollDataYGPU27(Parameter *para, vf::gpu::Communicator &comm, CudaMe
std::vector<ProcessNeighbor27> *sendProcessNeighborHost,
std::vector<ProcessNeighbor27> *recvProcessNeighborHost)
{
- cudaStream_t stream = para->getStreamManager()->getStream(CudaStreamIndex::Border);
+ cudaStream_t stream = para->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder);
///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
// copy Device to Host
for (unsigned int i = 0; i < (unsigned int)(para->getNumberOfProcessNeighborsY(level, "send")); i++)
diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
index 5f1393791..c3dd9e414 100644
--- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
+++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp
@@ -120,11 +120,11 @@ void GridGenerator::allocArrays_taggedFluidNodes() {
if(para->getParH(level)->numberOfTaggedFluidNodes[tag]>0)
para->getParH(level)->allocatedBulkFluidNodeTags.push_back(tag);
break;
- case CollisionTemplate::Border:
- this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesBorder(level), CollisionTemplate::Border, level);
- cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::Border, level);
- builder->getFluidNodeIndicesBorder(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::Border], level);
- cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::Border, level);
+ case CollisionTemplate::SubDomainBorder:
+ this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesBorder(level), CollisionTemplate::SubDomainBorder, level);
+ cudaMemoryManager->cudaAllocTaggedFluidNodeIndices(CollisionTemplate::SubDomainBorder, level);
+ builder->getFluidNodeIndicesBorder(para->getParH(level)->taggedFluidNodeIndices[CollisionTemplate::SubDomainBorder], level);
+ cudaMemoryManager->cudaCopyTaggedFluidNodeIndices(CollisionTemplate::SubDomainBorder, level);
break;
case CollisionTemplate::WriteMacroVars:
this->setNumberOfTaggedFluidNodes(builder->getNumberOfFluidNodesMacroVars(level), CollisionTemplate::WriteMacroVars, level);
@@ -157,7 +157,7 @@ void GridGenerator::allocArrays_taggedFluidNodes() {
VF_LOG_INFO("Number of tagged nodes on level {}:", level);
VF_LOG_INFO("Default: {}, Border: {}, WriteMacroVars: {}, ApplyBodyForce: {}, AllFeatures: {}",
para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default],
- para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::Border],
+ para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::SubDomainBorder],
para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::WriteMacroVars],
para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::ApplyBodyForce],
para->getParH(level)->numberOfTaggedFluidNodes[CollisionTemplate::AllFeatures] );
@@ -177,8 +177,8 @@ void GridGenerator::tagFluidNodeIndices(std::vector<uint> taggedFluidNodeIndices
builder->addFluidNodeIndicesAllFeatures( taggedFluidNodeIndices, level );
break;
case CollisionTemplate::Default:
- case CollisionTemplate::Border:
- throw std::runtime_error("Cannot tag fluid nodes as Default or Border!");
+ case CollisionTemplate::SubDomainBorder:
+ throw std::runtime_error("Cannot tag fluid nodes as Default or SubDomainBorder!");
default:
throw std::runtime_error("Tagging fluid nodes with invald tag!");
break;
diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
index 253960f96..8807c2230 100644
--- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
+++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp
@@ -530,7 +530,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborXIndex(int lev, unsigned int proc
void CudaMemoryManager::cudaCopyProcessNeighborXFsHD(int lev, unsigned int processNeighbor,
const unsigned int &memsizeFsRecv)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->recvProcessNeighborX[processNeighbor].f[0],
parameter->getParH(lev)->recvProcessNeighborX[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
@@ -540,12 +540,12 @@ void CudaMemoryManager::cudaCopyProcessNeighborXFsHD(int lev, unsigned int proce
parameter->getParH(lev)->recvProcessNeighborX[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
cudaMemcpyHostToDevice,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaCopyProcessNeighborXFsDH(int lev, unsigned int processNeighbor,
const unsigned int &memsizeFsSend)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->sendProcessNeighborX[processNeighbor].f[0],
parameter->getParD(lev)->sendProcessNeighborX[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
@@ -555,7 +555,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborXFsDH(int lev, unsigned int proce
parameter->getParD(lev)->sendProcessNeighborX[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
cudaMemcpyDeviceToHost,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaFreeProcessNeighborX(int lev, unsigned int processNeighbor)
{
@@ -600,7 +600,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborYIndex(int lev, unsigned int proc
}
void CudaMemoryManager::cudaCopyProcessNeighborYFsHD(int lev, unsigned int processNeighbor, const unsigned int &memsizeFsRecv)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->recvProcessNeighborY[processNeighbor].f[0],
parameter->getParH(lev)->recvProcessNeighborY[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
@@ -610,11 +610,11 @@ void CudaMemoryManager::cudaCopyProcessNeighborYFsHD(int lev, unsigned int proce
parameter->getParH(lev)->recvProcessNeighborY[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
cudaMemcpyHostToDevice,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaCopyProcessNeighborYFsDH(int lev, unsigned int processNeighbor, const unsigned int &memsizeFsSend)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->sendProcessNeighborY[processNeighbor].f[0],
parameter->getParD(lev)->sendProcessNeighborY[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
@@ -624,7 +624,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborYFsDH(int lev, unsigned int proce
parameter->getParD(lev)->sendProcessNeighborY[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
cudaMemcpyDeviceToHost,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaFreeProcessNeighborY(int lev, unsigned int processNeighbor)
{
@@ -670,7 +670,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborZIndex(int lev, unsigned int proc
void CudaMemoryManager::cudaCopyProcessNeighborZFsHD(int lev, unsigned int processNeighbor,
const unsigned int &memsizeFsRecv)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParD(lev)->recvProcessNeighborZ[processNeighbor].f[0],
parameter->getParH(lev)->recvProcessNeighborZ[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
@@ -680,12 +680,12 @@ void CudaMemoryManager::cudaCopyProcessNeighborZFsHD(int lev, unsigned int proce
parameter->getParH(lev)->recvProcessNeighborZ[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsRecv,
cudaMemcpyHostToDevice,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaCopyProcessNeighborZFsDH(int lev, unsigned int processNeighbor,
const unsigned int &memsizeFsSend)
{
- if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::Border))
+ if (!parameter->getStreamManager()->streamIsRegistered(CudaStreamIndex::SubDomainBorder))
checkCudaErrors( cudaMemcpy(parameter->getParH(lev)->sendProcessNeighborZ[processNeighbor].f[0],
parameter->getParD(lev)->sendProcessNeighborZ[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
@@ -695,7 +695,7 @@ void CudaMemoryManager::cudaCopyProcessNeighborZFsDH(int lev, unsigned int proce
parameter->getParD(lev)->sendProcessNeighborZ[processNeighbor].f[0],
parameter->getD3Qxx() * memsizeFsSend,
cudaMemcpyDeviceToHost,
- parameter->getStreamManager()->getStream(CudaStreamIndex::Border)));
+ parameter->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder)));
}
void CudaMemoryManager::cudaFreeProcessNeighborZ(int lev, unsigned int processNeighbor)
{
diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
index 54dd11142..8565d9660 100644
--- a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
+++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.cu
@@ -85,7 +85,7 @@ void CumulantK17<turbulenceModel>::runOnIndices( const unsigned int *indices, un
size_indices);
break;
- case CollisionTemplate::Border:
+ case CollisionTemplate::SubDomainBorder:
case CollisionTemplate::AllFeatures:
LB_Kernel_CumulantK17 < turbulenceModel, true, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>(
para->getParD(level)->omega,
diff --git a/src/gpu/VirtualFluids_GPU/LBM/LB.h b/src/gpu/VirtualFluids_GPU/LBM/LB.h
index e08544bdf..c07769d5c 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/LB.h
+++ b/src/gpu/VirtualFluids_GPU/LBM/LB.h
@@ -75,9 +75,9 @@ enum class CollisionTemplate {
//! - AllFeatures: collision \w write out macroscopic variables AND read and apply body force
AllFeatures,
//! - Border: collision on border nodes
- Border
+ SubDomainBorder
};
-constexpr std::initializer_list<CollisionTemplate> all_CollisionTemplate = { CollisionTemplate::Default, CollisionTemplate::WriteMacroVars, CollisionTemplate::ApplyBodyForce, CollisionTemplate::AllFeatures, CollisionTemplate::Border};
+constexpr std::initializer_list<CollisionTemplate> all_CollisionTemplate = { CollisionTemplate::Default, CollisionTemplate::WriteMacroVars, CollisionTemplate::ApplyBodyForce, CollisionTemplate::AllFeatures, CollisionTemplate::SubDomainBorder};
constexpr std::initializer_list<CollisionTemplate> bulk_CollisionTemplate = { CollisionTemplate::Default, CollisionTemplate::WriteMacroVars, CollisionTemplate::ApplyBodyForce, CollisionTemplate::AllFeatures};
struct InitCondition
diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
index b120dd205..84ab84ff9 100644
--- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
+++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp
@@ -139,7 +139,7 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa
//////////////////////////////////////////////////////////////////////////
// CUDA streams
if (para->getUseStreams()) {
- para->getStreamManager()->registerStream(CudaStreamIndex::Border);
+ para->getStreamManager()->registerStream(CudaStreamIndex::SubDomainBorder);
para->getStreamManager()->registerStream(CudaStreamIndex::Bulk);
para->getStreamManager()->launchStreams();
para->getStreamManager()->createCudaEvents();
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h b/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
index 06fc56247..5c59bcd3a 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/CudaStreamManager.h
@@ -39,7 +39,7 @@ enum class CudaStreamIndex
{
Legacy,
Bulk,
- Border,
+ SubDomainBorder,
Precursor,
ActuatorFarm
};
diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
index f83e50d52..86b7bc2a0 100644
--- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
+++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h
@@ -372,12 +372,12 @@ struct LBMSimulationParameter {
///////////////////////////////////////////////////////
std::map<CollisionTemplate, uint*> taggedFluidNodeIndices = {{CollisionTemplate::Default, nullptr},
- {CollisionTemplate::Border, nullptr},
+ {CollisionTemplate::SubDomainBorder,nullptr},
{CollisionTemplate::WriteMacroVars, nullptr},
{CollisionTemplate::ApplyBodyForce, nullptr},
{CollisionTemplate::AllFeatures, nullptr}};
std::map<CollisionTemplate, uint > numberOfTaggedFluidNodes = {{CollisionTemplate::Default, 0},
- {CollisionTemplate::Border, 0},
+ {CollisionTemplate::SubDomainBorder,0},
{CollisionTemplate::WriteMacroVars, 0},
{CollisionTemplate::ApplyBodyForce, 0},
{CollisionTemplate::AllFeatures, 0}};
--
GitLab