CollisisionStrategy.cpp

Go to the documentation of this file.

//=======================================================================================
// ____          ____    __    ______     __________   __      __       __        __
// \    \       |    |  |  |  |   _   \  |___    ___| |  |    |  |     /  \      |  |
//  \    \      |    |  |  |  |  |_)   |     |  |     |  |    |  |    /    \     |  |
//   \    \     |    |  |  |  |   _   /      |  |     |  |    |  |   /  /\  \    |  |
//    \    \    |    |  |  |  |  | \  \      |  |     |   \__/   |  /  ____  \   |  |____
//     \    \   |    |  |__|  |__|  \__\     |__|      \________/  /__/    \__\  |_______|
//      \    \  |    |   ________________________________________________________________
//       \    \ |    |  |  ______________________________________________________________|
//        \    \|    |  |  |         __          __     __     __     ______      _______
//         \         |  |  |_____   |  |        |  |   |  |   |  |   |   _  \    /  _____)
//          \        |  |   _____|  |  |        |  |   |  |   |  |   |  | \  \   \_______
//           \       |  |  |        |  |_____   |   \_/   |   |  |   |  |_/  /    _____  |
//            \ _____|  |__|        |________|   \_______/    |__|   |______/    (_______/
//
//  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.
//
//  SPDX-License-Identifier: GPL-3.0-or-later
//  SPDX-FileCopyrightText: Copyright © VirtualFluids Project contributors, see AUTHORS.md in root folder
//
//=======================================================================================
#include "CollisionStrategy.h"
#include "Cuda/CudaStreamManager.h"
#include "Parameter/Parameter.h"
#include <logger/Logger.h>
 
namespace vf::gpu {
 
std::function<void(UpdateGrid27 *updateGrid, Parameter *para, int level, unsigned int t)>
getFunctionForCollisionAndExchange(const bool useStreams, const int numberOfMpiProcesses,
                                   const bool kernelNeedsFluidNodeIndicesToRun)
{
    VF_LOG_INFO("Function used for collisionAndExchange: ");
 
    if (useStreams && numberOfMpiProcesses > 1 && kernelNeedsFluidNodeIndicesToRun) {
        VF_LOG_INFO("CollisionAndExchange_streams()");
        return CollisionAndExchange_streams();
 
    } else if (useStreams && !kernelNeedsFluidNodeIndicesToRun) {
        VF_LOG_INFO("Cuda Streams can only be used with kernels which run using fluidNodesIndices.");
 
    } else if (useStreams && numberOfMpiProcesses <= 1) {
        VF_LOG_INFO("Cuda Streams can only be used with multiple MPI processes.");
 
    } else if (!useStreams && kernelNeedsFluidNodeIndicesToRun) {
        VF_LOG_INFO("CollisionAndExchange_noStreams_indexKernel()");
        return CollisionAndExchange_noStreams_indexKernel();
 
    } else if (!useStreams && !kernelNeedsFluidNodeIndicesToRun) {
        VF_LOG_INFO("CollisionAndExchange_noStreams_oldKernel()");
        return CollisionAndExchange_noStreams_oldKernel();
    }
 
    throw std::runtime_error("Invalid Configuration for collision and exchange");
}
 
void CollisionAndExchange_noStreams_indexKernel::operator()(UpdateGrid27 *updateGrid, Parameter *para, int level,
                                                            unsigned int t)
{
    for( CollisionTemplate tag: para->getParH(level)->allocatedBulkFluidNodeTags )
    {
        updateGrid->collisionUsingIndices(  level, t, 
                                            para->getParD(level)->taggedFluidNodeIndices[tag],
                                            para->getParD(level)->numberOfTaggedFluidNodes[tag],
                                            tag,
                                            CudaStreamIndex::Legacy);
    }
 
    updateGrid->exchangeMultiGPU_noStreams_withPrepare(level, false);
}
 
void CollisionAndExchange_noStreams_oldKernel::operator()(UpdateGrid27 *updateGrid, Parameter *para, int level,
                                                          unsigned int t)
{
    updateGrid->collisionAllNodes(level, t);
 
    updateGrid->exchangeMultiGPU_noStreams_withPrepare(level, false);
}
 
void CollisionAndExchange_streams::operator()(UpdateGrid27 *updateGrid, Parameter *para, int level, unsigned int t)
{
    updateGrid->collisionUsingIndices(  level, t, 
                                        para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::SubDomainBorder],
                                        para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::SubDomainBorder], 
                                        CollisionTemplate::WriteMacroVars,  
                                        CudaStreamIndex::SubDomainBorder);
 
    updateGrid->prepareExchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
    if (para->getUseStreams())
        para->getStreamManager()->triggerStartBulkKernel(CudaStreamIndex::SubDomainBorder);
 
 
    para->getStreamManager()->waitOnStartBulkKernelEvent(CudaStreamIndex::Bulk);
    
    for( CollisionTemplate tag: para->getParH(level)->allocatedBulkFluidNodeTags )
    {
        updateGrid->collisionUsingIndices(  level, t, 
                                            para->getParD(level)->taggedFluidNodeIndices[tag],
                                            para->getParD(level)->numberOfTaggedFluidNodes[tag], 
                                            tag,
                                            CudaStreamIndex::Bulk);
    }
    updateGrid->exchangeMultiGPU(level, CudaStreamIndex::SubDomainBorder);
 
    cudaStreamSynchronize(para->getStreamManager()->getStream(CudaStreamIndex::SubDomainBorder));
    cudaStreamSynchronize(para->getStreamManager()->getStream(CudaStreamIndex::Bulk));
}
 
}