From e97c3e075fe62eb4bbcf48664a03580f31b82827 Mon Sep 17 00:00:00 2001 From: Martin Schoenherr <m.schoenherr@tu-braunschweig.de> Date: Fri, 17 Mar 2023 11:44:06 +0100 Subject: [PATCH] rename interpolation cell neighbor objects --- .../Calculation/RefinementStrategy.cpp | 28 ++++---- .../GridReaderFiles/GridReader.cpp | 4 +- .../GridReaderGenerator/GridGenerator.cpp | 4 +- .../InterpolationCellGrouper.cpp | 72 +++++++++---------- .../InterpolationCellGrouperTest.cpp | 36 +++++----- .../GPU/CudaMemoryManager.cpp | 48 ++++++------- .../GridScalingKernelManager.cpp | 8 +-- .../Output/InterfaceDebugWriter.hpp | 6 +- .../VirtualFluids_GPU/Output/OffsetWriter.hpp | 12 ++-- .../VirtualFluids_GPU/Parameter/Parameter.h | 8 +-- 10 files changed, 113 insertions(+), 113 deletions(-) diff --git a/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp b/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp index 61e6f37b2..f1c9bb8e5 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp +++ b/src/gpu/VirtualFluids_GPU/Calculation/RefinementStrategy.cpp @@ -42,7 +42,7 @@ 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)->fineToCoarseBorder, para->getParD(level)->neighborFC, CudaStreamIndex::SubDomainBorder); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarseBorder, para->getParD(level)->neighborFineToCoarse, 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 //! @@ -53,8 +53,8 @@ void RefinementAndExchange_streams_exchangeInterface::operator()(UpdateGrid27 *u //! 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)->fineToCoarseBulk, para->getParD(level)->neighborFCBulk, CudaStreamIndex::SubDomainBorder); - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBulk, para->getParD(level)->neighborCFBulk, CudaStreamIndex::SubDomainBorder); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarseBulk, para->getParD(level)->neighborFineToCoarseBulk, CudaStreamIndex::SubDomainBorder); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBulk, para->getParD(level)->neighborCoarseToFineBulk, CudaStreamIndex::SubDomainBorder); //! 4. exchange information between GPUs (only nodes which are part of the interpolation) //! @@ -62,7 +62,7 @@ void RefinementAndExchange_streams_exchangeInterface::operator()(UpdateGrid27 *u // 5. interpolation fine to coarse for nodes which are at the border of the gpus/processes //! - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBorder, para->getParD(level)->neighborCF, CudaStreamIndex::SubDomainBorder); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBorder, para->getParD(level)->neighborCoarseToFine, CudaStreamIndex::SubDomainBorder); cudaDeviceSynchronize(); } @@ -73,7 +73,7 @@ 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)->fineToCoarseBorder, para->getParD(level)->neighborFC, CudaStreamIndex::SubDomainBorder); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarseBorder, para->getParD(level)->neighborFineToCoarse, 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 //! @@ -84,8 +84,8 @@ void RefinementAndExchange_streams_exchangeAllNodes::operator()(UpdateGrid27 *up //! 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)->fineToCoarseBulk, para->getParD(level)->neighborFCBulk, CudaStreamIndex::SubDomainBorder); - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBulk, para->getParD(level)->neighborCFBulk, CudaStreamIndex::SubDomainBorder); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarseBulk, para->getParD(level)->neighborFineToCoarseBulk, CudaStreamIndex::SubDomainBorder); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBulk, para->getParD(level)->neighborCoarseToFineBulk, CudaStreamIndex::SubDomainBorder); //! 4. exchange information between GPUs (all nodes) //! @@ -93,7 +93,7 @@ void RefinementAndExchange_streams_exchangeAllNodes::operator()(UpdateGrid27 *up // 5. interpolation fine to coarse for nodes which are at the border of the gpus/processes //! - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBorder, para->getParD(level)->neighborCF, CudaStreamIndex::SubDomainBorder); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFineBorder, para->getParD(level)->neighborCoarseToFine, CudaStreamIndex::SubDomainBorder); cudaDeviceSynchronize(); } @@ -104,14 +104,14 @@ void RefinementAndExchange_noStreams_exchangeInterface::operator()(UpdateGrid27 //! //! 1. interpolation fine to coarse //! - updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFC, CudaStreamIndex::Legacy); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFineToCoarse, CudaStreamIndex::Legacy); //! 2. exchange information between GPUs (only nodes which are part of the interpolation) //! updateGrid->exchangeMultiGPU_noStreams_withPrepare(level, true); //! 3. interpolation coarse to fine - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCF, CudaStreamIndex::Legacy); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCoarseToFine, CudaStreamIndex::Legacy); } void RefinementAndExchange_noStreams_exchangeAllNodes::operator()(UpdateGrid27 *updateGrid, Parameter *para, int level) @@ -120,14 +120,14 @@ void RefinementAndExchange_noStreams_exchangeAllNodes::operator()(UpdateGrid27 * //! //! 1. interpolation fine to coarse //! - updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFC, CudaStreamIndex::Legacy); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFineToCoarse, CudaStreamIndex::Legacy); //! 2. exchange information between GPUs (all nodes) //! updateGrid->exchangeMultiGPU_noStreams_withPrepare(level, false); //! 3. interpolation coarse to fine - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCF, CudaStreamIndex::Legacy); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCoarseToFine, CudaStreamIndex::Legacy); } void Refinement_noExchange::operator()(UpdateGrid27 *updateGrid, Parameter *para, int level) @@ -136,7 +136,7 @@ void Refinement_noExchange::operator()(UpdateGrid27 *updateGrid, Parameter *para //! //! 1. interpolation fine to coarse //! - updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFC, CudaStreamIndex::Legacy); + updateGrid->fineToCoarse(level, ¶->getParD(level)->fineToCoarse, para->getParD(level)->neighborFineToCoarse, CudaStreamIndex::Legacy); //! 2. interpolation coarse to fine - updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCF, CudaStreamIndex::Legacy); + updateGrid->coarseToFine(level, ¶->getParD(level)->coarseToFine, para->getParD(level)->neighborCoarseToFine, CudaStreamIndex::Legacy); } diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp index 59473236f..647ab169b 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderFiles/GridReader.cpp @@ -180,8 +180,8 @@ void GridReader::allocArrays_OffsetScale() cudaMemoryManager->cudaAllocInterfaceOffFC(i); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// //init - obj_offCF->initArrayOffset(para->getParH(i)->neighborCF.x, para->getParH(i)->neighborCF.y, para->getParH(i)->neighborCF.z, i); - obj_offFC->initArrayOffset(para->getParH(i)->neighborFC.x, para->getParH(i)->neighborFC.y, para->getParH(i)->neighborFC.z, i); + obj_offCF->initArrayOffset(para->getParH(i)->neighborCoarseToFine.x, para->getParH(i)->neighborCoarseToFine.y, para->getParH(i)->neighborCoarseToFine.z, i); + obj_offFC->initArrayOffset(para->getParH(i)->neighborFineToCoarse.x, para->getParH(i)->neighborFineToCoarse.y, para->getParH(i)->neighborFineToCoarse.z, i); obj_scaleCFC->initScale(para->getParH(i)->coarseToFine.coarseCellIndices, i); obj_scaleCFF->initScale(para->getParH(i)->coarseToFine.fineCellIndices, i); obj_scaleFCC->initScale(para->getParH(i)->fineToCoarse.coarseCellIndices, i); diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp index 1139ad572..533382211 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/GridGenerator.cpp @@ -1194,8 +1194,8 @@ void GridGenerator::allocArrays_OffsetScale() cudaMemoryManager->cudaAllocInterfaceOffFC(level); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// //init - builder->getOffsetCF(para->getParH(level)->neighborCF.x, para->getParH(level)->neighborCF.y, para->getParH(level)->neighborCF.z, level); - builder->getOffsetFC(para->getParH(level)->neighborFC.x, para->getParH(level)->neighborFC.y, para->getParH(level)->neighborFC.z, level); + builder->getOffsetCF(para->getParH(level)->neighborCoarseToFine.x, para->getParH(level)->neighborCoarseToFine.y, para->getParH(level)->neighborCoarseToFine.z, level); + builder->getOffsetFC(para->getParH(level)->neighborFineToCoarse.x, para->getParH(level)->neighborFineToCoarse.y, para->getParH(level)->neighborFineToCoarse.z, level); builder->getGridInterfaceIndices(para->getParH(level)->coarseToFine.coarseCellIndices, para->getParH(level)->coarseToFine.fineCellIndices, para->getParH(level)->fineToCoarse.coarseCellIndices, para->getParH(level)->fineToCoarse.fineCellIndices, level); if (para->getUseStreams() || para->getNumprocs() > 1) { diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp index 1bc8e107b..f72536730 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouper.cpp @@ -20,9 +20,9 @@ void InterpolationCellGrouper::splitFineToCoarseIntoBorderAndBulk(uint level) co parDs[level]->fineToCoarseBulk.coarseCellIndices = parDs[level]->fineToCoarseBorder.coarseCellIndices + parDs[level]->fineToCoarseBorder.numberOfCells; parDs[level]->fineToCoarseBorder.fineCellIndices = parDs[level]->fineToCoarse.fineCellIndices; parDs[level]->fineToCoarseBulk.fineCellIndices = parDs[level]->fineToCoarseBorder.fineCellIndices + parDs[level]->fineToCoarseBorder.numberOfCells; - parDs[level]->neighborFCBulk.x = parDs[level]->neighborFC.x + parDs[level]->fineToCoarseBorder.numberOfCells; - parDs[level]->neighborFCBulk.y = parDs[level]->neighborFC.y + parDs[level]->fineToCoarseBorder.numberOfCells; - parDs[level]->neighborFCBulk.z = parDs[level]->neighborFC.z + parDs[level]->fineToCoarseBorder.numberOfCells; + parDs[level]->neighborFineToCoarseBulk.x = parDs[level]->neighborFineToCoarse.x + parDs[level]->fineToCoarseBorder.numberOfCells; + parDs[level]->neighborFineToCoarseBulk.y = parDs[level]->neighborFineToCoarse.y + parDs[level]->fineToCoarseBorder.numberOfCells; + parDs[level]->neighborFineToCoarseBulk.z = parDs[level]->neighborFineToCoarse.z + parDs[level]->fineToCoarseBorder.numberOfCells; } void InterpolationCellGrouper::reorderFineToCoarseIntoBorderAndBulk(uint level) const @@ -48,15 +48,15 @@ void InterpolationCellGrouper::reorderFineToCoarseIntoBorderAndBulk(uint level) if (grid->isSparseIndexInFluidNodeIndicesBorder(iCellFccAll[i])) { iCellFccBorderVector.push_back(iCellFccAll[i]); iCellFcfBorderVector.push_back(iCellFcfAll[i]); - xBorderVector.push_back(parHs[level]->neighborFC.x[i]); - yBorderVector.push_back(parHs[level]->neighborFC.y[i]); - zBorderVector.push_back(parHs[level]->neighborFC.z[i]); + xBorderVector.push_back(parHs[level]->neighborFineToCoarse.x[i]); + yBorderVector.push_back(parHs[level]->neighborFineToCoarse.y[i]); + zBorderVector.push_back(parHs[level]->neighborFineToCoarse.z[i]); } else { iCellFccBulkVector.push_back(iCellFccAll[i]); iCellFcfBulkVector.push_back(iCellFcfAll[i]); - xBulkVector.push_back(parHs[level]->neighborFC.x[i]); - yBulkVector.push_back(parHs[level]->neighborFC.y[i]); - zBulkVector.push_back(parHs[level]->neighborFC.z[i]); + xBulkVector.push_back(parHs[level]->neighborFineToCoarse.x[i]); + yBulkVector.push_back(parHs[level]->neighborFineToCoarse.y[i]); + zBulkVector.push_back(parHs[level]->neighborFineToCoarse.z[i]); } // set new sizes and pointers @@ -66,25 +66,25 @@ void InterpolationCellGrouper::reorderFineToCoarseIntoBorderAndBulk(uint level) parHs[level]->fineToCoarseBulk.numberOfCells = (uint)iCellFccBulkVector.size(); parHs[level]->fineToCoarseBulk.coarseCellIndices = iCellFccAll + parHs[level]->fineToCoarseBorder.numberOfCells; parHs[level]->fineToCoarseBulk.fineCellIndices = iCellFcfAll + parHs[level]->fineToCoarseBorder.numberOfCells; - parHs[level]->neighborFCBulk.x = parHs[level]->neighborFC.x + parHs[level]->fineToCoarseBorder.numberOfCells; - parHs[level]->neighborFCBulk.y = parHs[level]->neighborFC.y + parHs[level]->fineToCoarseBorder.numberOfCells; - parHs[level]->neighborFCBulk.z = parHs[level]->neighborFC.z + parHs[level]->fineToCoarseBorder.numberOfCells; + parHs[level]->neighborFineToCoarseBulk.x = parHs[level]->neighborFineToCoarse.x + parHs[level]->fineToCoarseBorder.numberOfCells; + parHs[level]->neighborFineToCoarseBulk.y = parHs[level]->neighborFineToCoarse.y + parHs[level]->fineToCoarseBorder.numberOfCells; + parHs[level]->neighborFineToCoarseBulk.z = parHs[level]->neighborFineToCoarse.z + parHs[level]->fineToCoarseBorder.numberOfCells; // copy the created vectors to the memory addresses of the old arrays // this is inefficient :( for (uint i = 0; i < (uint)iCellFccBorderVector.size(); i++) { iCellFccAll[i] = iCellFccBorderVector[i]; iCellFcfAll[i] = iCellFcfBorderVector[i]; - parHs[level]->neighborFC.x[i] = xBorderVector[i]; - parHs[level]->neighborFC.y[i] = yBorderVector[i]; - parHs[level]->neighborFC.z[i] = zBorderVector[i]; + parHs[level]->neighborFineToCoarse.x[i] = xBorderVector[i]; + parHs[level]->neighborFineToCoarse.y[i] = yBorderVector[i]; + parHs[level]->neighborFineToCoarse.z[i] = zBorderVector[i]; } for (uint i = 0; i < (uint)iCellFccBulkVector.size(); i++) { parHs[level]->fineToCoarseBulk.coarseCellIndices[i] = iCellFccBulkVector[i]; parHs[level]->fineToCoarseBulk.fineCellIndices[i] = iCellFcfBulkVector[i]; - parHs[level]->neighborFCBulk.x[i] = xBulkVector[i]; - parHs[level]->neighborFCBulk.y[i] = yBulkVector[i]; - parHs[level]->neighborFCBulk.z[i] = zBulkVector[i]; + parHs[level]->neighborFineToCoarseBulk.x[i] = xBulkVector[i]; + parHs[level]->neighborFineToCoarseBulk.y[i] = yBulkVector[i]; + parHs[level]->neighborFineToCoarseBulk.z[i] = zBulkVector[i]; } } @@ -98,9 +98,9 @@ void InterpolationCellGrouper::splitCoarseToFineIntoBorderAndBulk(uint level) co parDs[level]->coarseToFineBulk.coarseCellIndices = parDs[level]->coarseToFineBorder.coarseCellIndices + parDs[level]->coarseToFineBorder.numberOfCells; parDs[level]->coarseToFineBorder.fineCellIndices = parDs[level]->coarseToFine.fineCellIndices; parDs[level]->coarseToFineBulk.fineCellIndices = parDs[level]->coarseToFineBorder.fineCellIndices + parDs[level]->coarseToFineBorder.numberOfCells; - parDs[level]->neighborCFBulk.x = parDs[level]->neighborCF.x + parDs[level]->coarseToFineBorder.numberOfCells; - parDs[level]->neighborCFBulk.y = parDs[level]->neighborCF.y + parDs[level]->coarseToFineBorder.numberOfCells; - parDs[level]->neighborCFBulk.z = parDs[level]->neighborCF.z + parDs[level]->coarseToFineBorder.numberOfCells; + parDs[level]->neighborCoarseToFineBulk.x = parDs[level]->neighborCoarseToFine.x + parDs[level]->coarseToFineBorder.numberOfCells; + parDs[level]->neighborCoarseToFineBulk.y = parDs[level]->neighborCoarseToFine.y + parDs[level]->coarseToFineBorder.numberOfCells; + parDs[level]->neighborCoarseToFineBulk.z = parDs[level]->neighborCoarseToFine.z + parDs[level]->coarseToFineBorder.numberOfCells; } void InterpolationCellGrouper::reorderCoarseToFineIntoBorderAndBulk(uint level) const @@ -140,15 +140,15 @@ void InterpolationCellGrouper::reorderCoarseToFineIntoBorderAndBulk(uint level) iCellCfcBorderVector.push_back(iCellCfcAll[i]); iCellCffBorderVector.push_back(iCellCffAll[i]); - xBorderVector.push_back(parHs[level]->neighborCF.x[i]); - yBorderVector.push_back(parHs[level]->neighborCF.y[i]); - zBorderVector.push_back(parHs[level]->neighborCF.z[i]); + xBorderVector.push_back(parHs[level]->neighborCoarseToFine.x[i]); + yBorderVector.push_back(parHs[level]->neighborCoarseToFine.y[i]); + zBorderVector.push_back(parHs[level]->neighborCoarseToFine.z[i]); } else { iCellCfcBulkVector.push_back(iCellCfcAll[i]); iCellCffBulkVector.push_back(iCellCffAll[i]); - xBulkVector.push_back(parHs[level]->neighborCF.x[i]); - yBulkVector.push_back(parHs[level]->neighborCF.y[i]); - zBulkVector.push_back(parHs[level]->neighborCF.z[i]); + xBulkVector.push_back(parHs[level]->neighborCoarseToFine.x[i]); + yBulkVector.push_back(parHs[level]->neighborCoarseToFine.y[i]); + zBulkVector.push_back(parHs[level]->neighborCoarseToFine.z[i]); } } @@ -159,24 +159,24 @@ void InterpolationCellGrouper::reorderCoarseToFineIntoBorderAndBulk(uint level) parHs[level]->coarseToFineBulk.numberOfCells = (uint)iCellCfcBulkVector.size(); parHs[level]->coarseToFineBulk.coarseCellIndices = parHs[level]->coarseToFine.coarseCellIndices + parHs[level]->coarseToFineBorder.numberOfCells; parHs[level]->coarseToFineBulk.fineCellIndices = parHs[level]->coarseToFine.fineCellIndices + parHs[level]->coarseToFineBorder.numberOfCells; - parHs[level]->neighborCFBulk.x = parHs[level]->neighborCF.x + parHs[level]->coarseToFineBorder.numberOfCells; - parHs[level]->neighborCFBulk.y = parHs[level]->neighborCF.y + parHs[level]->coarseToFineBorder.numberOfCells; - parHs[level]->neighborCFBulk.z = parHs[level]->neighborCF.z + parHs[level]->coarseToFineBorder.numberOfCells; + parHs[level]->neighborCoarseToFineBulk.x = parHs[level]->neighborCoarseToFine.x + parHs[level]->coarseToFineBorder.numberOfCells; + parHs[level]->neighborCoarseToFineBulk.y = parHs[level]->neighborCoarseToFine.y + parHs[level]->coarseToFineBorder.numberOfCells; + parHs[level]->neighborCoarseToFineBulk.z = parHs[level]->neighborCoarseToFine.z + parHs[level]->coarseToFineBorder.numberOfCells; // copy the created vectors to the memory addresses of the old arrays // this is inefficient :( for (uint i = 0; i < (uint)iCellCfcBorderVector.size(); i++) { parHs[level]->coarseToFineBorder.coarseCellIndices[i] = iCellCfcBorderVector[i]; parHs[level]->coarseToFineBorder.fineCellIndices[i] = iCellCffBorderVector[i]; - parHs[level]->neighborCF.x[i] = xBorderVector[i]; - parHs[level]->neighborCF.y[i] = yBorderVector[i]; - parHs[level]->neighborCF.z[i] = zBorderVector[i]; + parHs[level]->neighborCoarseToFine.x[i] = xBorderVector[i]; + parHs[level]->neighborCoarseToFine.y[i] = yBorderVector[i]; + parHs[level]->neighborCoarseToFine.z[i] = zBorderVector[i]; } for (uint i = 0; i < (uint)iCellCfcBulkVector.size(); i++) { parHs[level]->coarseToFineBulk.coarseCellIndices[i] = iCellCfcBulkVector[i]; parHs[level]->coarseToFineBulk.fineCellIndices[i] = iCellCffBulkVector[i]; - parHs[level]->neighborCFBulk.x[i] = xBulkVector[i]; - parHs[level]->neighborCFBulk.y[i] = yBulkVector[i]; - parHs[level]->neighborCFBulk.z[i] = zBulkVector[i]; + parHs[level]->neighborCoarseToFineBulk.x[i] = xBulkVector[i]; + parHs[level]->neighborCoarseToFineBulk.y[i] = yBulkVector[i]; + parHs[level]->neighborCoarseToFineBulk.z[i] = zBulkVector[i]; } } diff --git a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp index 83b63f73e..6a8b28c09 100644 --- a/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp +++ b/src/gpu/VirtualFluids_GPU/DataStructureInitializer/GridReaderGenerator/InterpolationCellGrouperTest.cpp @@ -111,9 +111,9 @@ private: para->getParH(cf.level)->neighborY = cf.neighborY; para->getParH(cf.level)->neighborZ = cf.neighborZ; para->getParH(cf.level)->coarseToFine.numberOfCells = cf.sizeOfICellCf; - para->getParH(cf.level)->neighborCF.x = &(cf.offsetCFx.front()); - para->getParH(cf.level)->neighborCF.y = &(cf.offsetCFy.front()); - para->getParH(cf.level)->neighborCF.z = &(cf.offsetCFz.front()); + para->getParH(cf.level)->neighborCoarseToFine.x = &(cf.offsetCFx.front()); + para->getParH(cf.level)->neighborCoarseToFine.y = &(cf.offsetCFy.front()); + para->getParH(cf.level)->neighborCoarseToFine.z = &(cf.offsetCFz.front()); return std::make_unique<InterpolationCellGrouper>(para->getParHallLevels(), para->getParDallLevels(), builder); }; @@ -151,12 +151,12 @@ TEST_F(InterpolationCellGrouperTest_IndicesCFBorderBulkTest, splitCoarseToFineIn << "coarseToFineBulk.ICellCFF does not match the expected bulk vector"; // check offset cells - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCF.x, cf.offsetCFx_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCFBulk.x, cf.offsetCFx_Bulk_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCF.y, cf.offsetCFy_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCFBulk.y, cf.offsetCFy_Bulk_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCF.z, cf.offsetCFz_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCFBulk.z, cf.offsetCFz_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFine.x, cf.offsetCFx_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFineBulk.x, cf.offsetCFx_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFine.y, cf.offsetCFy_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFineBulk.y, cf.offsetCFy_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFine.z, cf.offsetCFz_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(cf.level)->neighborCoarseToFineBulk.z, cf.offsetCFz_Bulk_expected)); } struct FCBorderBulk { @@ -202,9 +202,9 @@ private: para->getParH(fc.level)->fineToCoarse.coarseCellIndices = &(fc.iCellFCC.front()); para->getParH(fc.level)->fineToCoarse.fineCellIndices = &(fc.iCellFCF.front()); para->getParH(fc.level)->fineToCoarse.numberOfCells = fc.sizeOfICellFC; - para->getParH(fc.level)->neighborFC.x = &(fc.offsetFCx.front()); - para->getParH(fc.level)->neighborFC.y = &(fc.offsetFCy.front()); - para->getParH(fc.level)->neighborFC.z = &(fc.offsetFCz.front()); + para->getParH(fc.level)->neighborFineToCoarse.x = &(fc.offsetFCx.front()); + para->getParH(fc.level)->neighborFineToCoarse.y = &(fc.offsetFCy.front()); + para->getParH(fc.level)->neighborFineToCoarse.z = &(fc.offsetFCz.front()); return std::make_unique<InterpolationCellGrouper>(para->getParHallLevels(), para->getParDallLevels(), builder); }; @@ -242,10 +242,10 @@ TEST_F(InterpolationCellGrouperTest_IndicesFCBorderBulkTest, splitFineToCoarseIn << "fineToCoarseBulk.ICellFCF does not match the expected bulk vector"; // check offset cells - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFC.x, fc.offsetFCx_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFCBulk.x, fc.offsetFCx_Bulk_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFC.y, fc.offsetFCy_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFCBulk.y, fc.offsetFCy_Bulk_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFC.z, fc.offsetFCz_Border_expected)); - EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFCBulk.z, fc.offsetFCz_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarse.x, fc.offsetFCx_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarseBulk.x, fc.offsetFCx_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarse.y, fc.offsetFCy_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarseBulk.y, fc.offsetFCy_Bulk_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarse.z, fc.offsetFCz_Border_expected)); + EXPECT_TRUE(vectorsAreEqual(para->getParH(fc.level)->neighborFineToCoarseBulk.z, fc.offsetFCz_Bulk_expected)); } diff --git a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp index 968e11cb0..e9d40ddf7 100644 --- a/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp +++ b/src/gpu/VirtualFluids_GPU/GPU/CudaMemoryManager.cpp @@ -1170,14 +1170,14 @@ void CudaMemoryManager::cudaAllocInterfaceOffCF(int lev) uint mem_size_kCF_off = sizeof(real) * parameter->getParH(lev)->coarseToFine.numberOfCells; //Host - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCF.x), mem_size_kCF_off )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCF.y), mem_size_kCF_off )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCF.z), mem_size_kCF_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCoarseToFine.x), mem_size_kCF_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCoarseToFine.y), mem_size_kCF_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborCoarseToFine.z), mem_size_kCF_off )); getLastCudaError("Allocate host memory"); //Device - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCF.x), mem_size_kCF_off )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCF.y), mem_size_kCF_off )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCF.z), mem_size_kCF_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCoarseToFine.x), mem_size_kCF_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCoarseToFine.y), mem_size_kCF_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborCoarseToFine.z), mem_size_kCF_off )); getLastCudaError("Allocate device memory"); ////////////////////////////////////////////////////////////////////////// double tmp = 3. * (double)mem_size_kCF_off; @@ -1187,16 +1187,16 @@ void CudaMemoryManager::cudaCopyInterfaceOffCF(int lev) { uint mem_size_kCF_off = sizeof(real) * parameter->getParH(lev)->coarseToFine.numberOfCells; - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCF.x, parameter->getParH(lev)->neighborCF.x, mem_size_kCF_off, cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCF.y, parameter->getParH(lev)->neighborCF.y, mem_size_kCF_off, cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCF.z, parameter->getParH(lev)->neighborCF.z, mem_size_kCF_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCoarseToFine.x, parameter->getParH(lev)->neighborCoarseToFine.x, mem_size_kCF_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCoarseToFine.y, parameter->getParH(lev)->neighborCoarseToFine.y, mem_size_kCF_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborCoarseToFine.z, parameter->getParH(lev)->neighborCoarseToFine.z, mem_size_kCF_off, cudaMemcpyHostToDevice)); getLastCudaError("Copy host memory to device"); } void CudaMemoryManager::cudaFreeInterfaceOffCF(int lev) { - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCF.x)); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCF.y)); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCF.z)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCoarseToFine.x)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCoarseToFine.y)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborCoarseToFine.z)); } //Interface Offset FC void CudaMemoryManager::cudaAllocInterfaceOffFC(int lev) @@ -1204,14 +1204,14 @@ void CudaMemoryManager::cudaAllocInterfaceOffFC(int lev) uint mem_size_kFC_off = sizeof(real) * parameter->getParH(lev)->fineToCoarse.numberOfCells; //Host - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFC.x), mem_size_kFC_off )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFC.y), mem_size_kFC_off )); - checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFC.z), mem_size_kFC_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFineToCoarse.x), mem_size_kFC_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFineToCoarse.y), mem_size_kFC_off )); + checkCudaErrors( cudaMallocHost((void**) &(parameter->getParH(lev)->neighborFineToCoarse.z), mem_size_kFC_off )); getLastCudaError("Allocate host memory"); //Device - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFC.x), mem_size_kFC_off )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFC.y), mem_size_kFC_off )); - checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFC.z), mem_size_kFC_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFineToCoarse.x), mem_size_kFC_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFineToCoarse.y), mem_size_kFC_off )); + checkCudaErrors( cudaMalloc((void**) &(parameter->getParD(lev)->neighborFineToCoarse.z), mem_size_kFC_off )); getLastCudaError("Allocate device memory"); ////////////////////////////////////////////////////////////////////////// double tmp = 3. * (double)mem_size_kFC_off; @@ -1221,16 +1221,16 @@ void CudaMemoryManager::cudaCopyInterfaceOffFC(int lev) { uint mem_size_kFC_off = sizeof(real) * parameter->getParH(lev)->fineToCoarse.numberOfCells; - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFC.x, parameter->getParH(lev)->neighborFC.x, mem_size_kFC_off, cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFC.y, parameter->getParH(lev)->neighborFC.y, mem_size_kFC_off, cudaMemcpyHostToDevice)); - checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFC.z, parameter->getParH(lev)->neighborFC.z, mem_size_kFC_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFineToCoarse.x, parameter->getParH(lev)->neighborFineToCoarse.x, mem_size_kFC_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFineToCoarse.y, parameter->getParH(lev)->neighborFineToCoarse.y, mem_size_kFC_off, cudaMemcpyHostToDevice)); + checkCudaErrors(cudaMemcpy(parameter->getParD(lev)->neighborFineToCoarse.z, parameter->getParH(lev)->neighborFineToCoarse.z, mem_size_kFC_off, cudaMemcpyHostToDevice)); getLastCudaError("Copy host memory to device"); } void CudaMemoryManager::cudaFreeInterfaceOffFC(int lev) { - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFC.x)); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFC.y)); - checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFC.z)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFineToCoarse.x)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFineToCoarse.y)); + checkCudaErrors( cudaFreeHost(parameter->getParH(lev)->neighborFineToCoarse.z)); } //Inlet diff --git a/src/gpu/VirtualFluids_GPU/KernelManager/GridScalingKernelManager.cpp b/src/gpu/VirtualFluids_GPU/KernelManager/GridScalingKernelManager.cpp index 7f08a14e6..bf025694b 100644 --- a/src/gpu/VirtualFluids_GPU/KernelManager/GridScalingKernelManager.cpp +++ b/src/gpu/VirtualFluids_GPU/KernelManager/GridScalingKernelManager.cpp @@ -300,7 +300,7 @@ void GridScalingKernelManager::runFineToCoarseKernelAD(const int level) const para->getParD(level)->viscosity, para->getParD(level)->diffusivity, para->getParD(level)->numberofthreads, - para->getParD(level)->neighborFC); + para->getParD(level)->neighborFineToCoarse); } else if (para->getDiffMod() == 27) { @@ -324,7 +324,7 @@ void GridScalingKernelManager::runFineToCoarseKernelAD(const int level) const para->getParD(level)->viscosity, para->getParD(level)->diffusivity, para->getParD(level)->numberofthreads, - para->getParD(level)->neighborFC); + para->getParD(level)->neighborFineToCoarse); } } @@ -569,7 +569,7 @@ void GridScalingKernelManager::runCoarseToFineKernelAD(const int level) const para->getParD(level)->viscosity, para->getParD(level+1)->diffusivity, para->getParD(level)->numberofthreads, - para->getParD(level)->neighborCF); + para->getParD(level)->neighborCoarseToFine); } else if (para->getDiffMod() == 27) { @@ -593,6 +593,6 @@ void GridScalingKernelManager::runCoarseToFineKernelAD(const int level) const para->getParD(level)->viscosity, para->getParD(level+1)->diffusivity, para->getParD(level)->numberofthreads, - para->getParD(level)->neighborCF); + para->getParD(level)->neighborCoarseToFine); } } diff --git a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.hpp b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.hpp index b807e3f36..30213c8b3 100644 --- a/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.hpp +++ b/src/gpu/VirtualFluids_GPU/Output/InterfaceDebugWriter.hpp @@ -136,9 +136,9 @@ void writeInterfaceLinesDebugOff(Parameter *para) int nodeCount = 0; for (int level = 0; level < para->getMaxLevel(); level++) { for (unsigned int u = 0; u < para->getParH(level)->coarseToFine.numberOfCells; u++) { - double xoff = para->getParH(level)->neighborCF.x[u]; - double yoff = para->getParH(level)->neighborCF.y[u]; - double zoff = para->getParH(level)->neighborCF.z[u]; + double xoff = para->getParH(level)->neighborCoarseToFine.x[u]; + double yoff = para->getParH(level)->neighborCoarseToFine.y[u]; + double zoff = para->getParH(level)->neighborCoarseToFine.z[u]; int posFine = para->getParH(level)->coarseToFine.fineCellIndices[u]; diff --git a/src/gpu/VirtualFluids_GPU/Output/OffsetWriter.hpp b/src/gpu/VirtualFluids_GPU/Output/OffsetWriter.hpp index 081b062d8..fb04951db 100644 --- a/src/gpu/VirtualFluids_GPU/Output/OffsetWriter.hpp +++ b/src/gpu/VirtualFluids_GPU/Output/OffsetWriter.hpp @@ -25,9 +25,9 @@ public: out.writeLine(); for (unsigned int u = 0; u < para->getParH(level)->coarseToFine.numberOfCells; u++) { - out.writeDouble(para->getParH(level)->neighborCF.x[u]); - out.writeDouble(para->getParH(level)->neighborCF.y[u]); - out.writeDouble(para->getParH(level)->neighborCF.z[u]); + out.writeDouble(para->getParH(level)->neighborCoarseToFine.x[u]); + out.writeDouble(para->getParH(level)->neighborCoarseToFine.y[u]); + out.writeDouble(para->getParH(level)->neighborCoarseToFine.z[u]); } out.writeLine(); } //end levelloop @@ -40,9 +40,9 @@ public: out.writeLine(); for (unsigned int u = 0; u < para->getParH(level)->fineToCoarse.numberOfCells; u++) { - out.writeDouble(para->getParH(level)->neighborFC.x[u]); - out.writeDouble(para->getParH(level)->neighborFC.y[u]); - out.writeDouble(para->getParH(level)->neighborFC.z[u]); + out.writeDouble(para->getParH(level)->neighborFineToCoarse.x[u]); + out.writeDouble(para->getParH(level)->neighborFineToCoarse.y[u]); + out.writeDouble(para->getParH(level)->neighborFineToCoarse.z[u]); } out.writeLine(); } //end levelloop diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h index f4c5da995..b4b49e153 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h @@ -162,10 +162,10 @@ struct LBMSimulationParameter { InterpolationCell coarseToFineBulk; ////////////////////////////////////////////////////////////////////////// //! \brief stores location of neighboring cell (necessary for refinement into the wall) - InterpolationCellNeighbor neighborCF; - InterpolationCellNeighbor neighborCFBulk; - InterpolationCellNeighbor neighborFC; - InterpolationCellNeighbor neighborFCBulk; + InterpolationCellNeighbor neighborCoarseToFine; + InterpolationCellNeighbor neighborCoarseToFineBulk; + InterpolationCellNeighbor neighborFineToCoarse; + InterpolationCellNeighbor neighborFineToCoarseBulk; ////////////////////////////////////////////////////////////////////////// -- GitLab