From d3fda3960c3b7c2a7089e174863606f66edcc369 Mon Sep 17 00:00:00 2001
From: Anna Wellmann <a.wellmann@tu-bs.de>
Date: Thu, 31 Mar 2022 13:05:31 +0200
Subject: [PATCH] Add refinementAndExchange_streams_completeExchange() to
UpdateGrid
---
.../Calculation/UpdateGrid27.cpp | 50 ++++++++++++++++---
.../Calculation/UpdateGrid27.h | 5 +-
2 files changed, 47 insertions(+), 8 deletions(-)
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
index c47c14521..cd5355317 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp
@@ -48,7 +48,7 @@ void UpdateGrid27::updateGrid(int level, unsigned int t)
void UpdateGrid27::refinementAndExchange_noRefinementAndExchange(int level) {}
-void UpdateGrid27::refinementAndExchange_streams(int level)
+void UpdateGrid27::refinementAndExchange_streams_onlyExchangeInterface(int level)
{
int borderStreamIndex = para->getStreamManager()->getBorderStreamIndex();
int bulkStreamIndex = para->getStreamManager()->getBulkStreamIndex();
@@ -82,6 +82,40 @@ void UpdateGrid27::refinementAndExchange_streams(int level)
cudaDeviceSynchronize();
}
+void UpdateGrid27::refinementAndExchange_streams_completeExchange(int level)
+{
+ int borderStreamIndex = para->getStreamManager()->getBorderStreamIndex();
+ int bulkStreamIndex = para->getStreamManager()->getBulkStreamIndex();
+
+ // fine to coarse border
+ fineToCoarseWithStream(para.get(), level, para->getParD(level)->intFCBorder.ICellFCC,
+ para->getParD(level)->intFCBorder.ICellFCF, para->getParD(level)->intFCBorder.kFC,
+ borderStreamIndex);
+
+ // prepare exchange and trigger bulk kernel when finished
+ prepareExchangeMultiGPU(para.get(), level, borderStreamIndex);
+ if (para->getUseStreams())
+ para->getStreamManager()->triggerStartBulkKernel(borderStreamIndex);
+
+ // launch bulk kernels (f to c and c to f)
+ para->getStreamManager()->waitOnStartBulkKernelEvent(bulkStreamIndex);
+ fineToCoarseWithStream(para.get(), level, para->getParD(level)->intFCBulk.ICellFCC,
+ para->getParD(level)->intFCBulk.ICellFCF, para->getParD(level)->intFCBulk.kFC,
+ bulkStreamIndex);
+ coarseToFineWithStream(para.get(), level, para->getParD(level)->intCFBulk.ICellCFC,
+ para->getParD(level)->intCFBulk.ICellCFF, para->getParD(level)->intCFBulk.kCF, para->getParD(level)->offCFBulk,
+ bulkStreamIndex);
+
+ // exchange
+ exchangeMultiGPU(para.get(), comm, cudaManager.get(), level, borderStreamIndex);
+
+ // coarse to fine border
+ coarseToFineWithStream(para.get(), level, para->getParD(level)->intCFBorder.ICellCFC,
+ para->getParD(level)->intCFBorder.ICellCFF, para->getParD(level)->intCFBorder.kCF, para->getParD(level)->offCF,
+ borderStreamIndex);
+ cudaDeviceSynchronize();
+}
+
void UpdateGrid27::refinementAndExchange_noStreams_onlyExchangeInterface(int level)
{
fineToCoarse(para.get(), level);
@@ -1533,14 +1567,18 @@ void UpdateGrid27::chooseFunctionForRefinementAndExchange()
this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_noExchange;
std::cout << "refinementAndExchange_noExchange()" << std::endl;
- } else if (para->getUseStreams() && para->getNumprocs() > 1 && para->useReducedCommunicationAfterFtoC) {
- this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_streams;
- std::cout << "refinementAndExchange_streams()" << std::endl;
+ } else if (para->getNumprocs() > 1 && para->getUseStreams() && para->useReducedCommunicationAfterFtoC) {
+ this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_streams_onlyExchangeInterface;
+ std::cout << "refinementAndExchange_streams_onlyExchangeInterface()" << std::endl;
+
+ } else if(para->getNumprocs() > 1 && para->getUseStreams() && !para->useReducedCommunicationAfterFtoC){
+ this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_streams_completeExchange;
+ std::cout << "refinementAndExchange_streams_completeExchange()" << std::endl;
- } else if (para->getNumprocs() > 1 && para->useReducedCommunicationAfterFtoC) {
+ } else if (para->getNumprocs() > 1 && !para->getUseStreams() && para->useReducedCommunicationAfterFtoC) {
this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_noStreams_onlyExchangeInterface;
std::cout << "refinementAndExchange_noStreams_onlyExchangeInterface()" << std::endl;
-
+
} else {
this->refinementAndExchange = &UpdateGrid27::refinementAndExchange_noStreams_completeExchange;
std::cout << "refinementAndExchange_noStreams_completeExchange()" << std::endl;
diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
index d55e4ee72..086b39fcb 100644
--- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
+++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h
@@ -32,10 +32,11 @@ private:
void collisionAndExchange_streams(int level, unsigned int t);
// functions for refinement and exchange
- void refinementAndExchange_noRefinementAndExchange(int level);
- void refinementAndExchange_streams(int level);
+ void refinementAndExchange_streams_onlyExchangeInterface(int level);
+ void refinementAndExchange_streams_completeExchange(int level);
void refinementAndExchange_noStreams_onlyExchangeInterface(int level);
void refinementAndExchange_noStreams_completeExchange(int level);
+ void refinementAndExchange_noRefinementAndExchange(int level);
void refinementAndExchange_noExchange(int level);
--
GitLab