From 798ed020e5e11e7b379e981aa147d4aac83da07b Mon Sep 17 00:00:00 2001 From: Steffen Date: Fri, 13 Feb 2026 16:16:03 -0500 Subject: [PATCH 1/2] add mpi cuda --- CMakeLists.txt | 2 +- README.md | 6 +- include/solvers/snsolver_hpc_cuda.hpp | 2 +- src/solvers/snsolver_hpc.cu | 59 ++++++++++++++----- .../install_kitrt_singularity_cuda.sh | 2 +- 5 files changed, 49 insertions(+), 22 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 894a3d99..32bab039 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -9,7 +9,7 @@ option( BUILD_UNITY "enables unity build for faster compile times" ON ) option( BUILD_CODE_COV "enables compiler option required for code coverage analysis" OFF ) option( BUILD_ML "enables build with tensorflow backend access" OFF ) option( BUILD_MPI "enables build with MPI access" OFF ) -option( BUILD_CUDA_HPC "enables CUDA backend for SN HPC solver (single GPU)" OFF ) +option( BUILD_CUDA_HPC "enables CUDA backend for SN HPC solver (MPI rank to GPU mapping)" OFF ) ################################################# diff --git a/README.md b/README.md index 987ab671..33f62ebd 100644 --- a/README.md +++ b/README.md @@ -142,7 +142,7 @@ singularity exec tools/singularity/kit_rt_MPI.sif \ mpirun -np 4 ./build_singularity_mpi/KiT-RT tests/input/validation_tests/SN_solver_hpc/lattice_hpc_200_cpu_order2.cfg ``` -### 3. CPU + single GPU (OpenMP + CUDA) +### 3. CPU + CUDA (single or multi-GPU via MPI) #### 3a) Singularity installation ```bash @@ -152,11 +152,11 @@ cd ../.. mkdir -p build_singularity_cuda cd build_singularity_cuda singularity exec --nv ../tools/singularity/kit_rt_MPI_cuda.sif \ - cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_MPI=OFF -DBUILD_CUDA_HPC=ON -DBUILD_ML=OFF .. + cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_MPI=ON -DBUILD_CUDA_HPC=ON -DBUILD_ML=OFF .. singularity exec --nv ../tools/singularity/kit_rt_MPI_cuda.sif make -j cd .. singularity exec --nv tools/singularity/kit_rt_MPI_cuda.sif \ - ./build_singularity_cuda/KiT-RT tests/input/validation_tests/SN_solver_hpc/lattice_hpc_200_cuda_order2.cfg + mpirun -np 2 ./build_singularity_cuda/KiT-RT tests/input/validation_tests/SN_solver_hpc/lattice_hpc_200_cuda_order2.cfg ``` When compiled with `-DBUILD_CUDA_HPC=ON`, HPC runs use the CUDA backend if a GPU is visible, and fall back to CPU if no GPU is detected. diff --git a/include/solvers/snsolver_hpc_cuda.hpp b/include/solvers/snsolver_hpc_cuda.hpp index 5a6f73f0..651d729e 100644 --- a/include/solvers/snsolver_hpc_cuda.hpp +++ b/include/solvers/snsolver_hpc_cuda.hpp @@ -151,7 +151,7 @@ class SNSolverHPCCUDA std::vector _historyOutputFields; /*!< @brief Solver Output: dimensions (FieldID). */ std::vector _historyOutputFieldNames; /*!< @brief Names of the outputFields: dimensions (FieldID) */ - // CUDA backend (single GPU for first version) + // CUDA backend state bool _cudaInitialized; int _cudaDeviceId; DeviceBuffers* _device; diff --git a/src/solvers/snsolver_hpc.cu b/src/solvers/snsolver_hpc.cu index e016c797..c6facfca 100644 --- a/src/solvers/snsolver_hpc.cu +++ b/src/solvers/snsolver_hpc.cu @@ -369,21 +369,14 @@ SNSolverHPCCUDA::SNSolverHPCCUDA( Config* settings ) { ErrorMessages::Error( "The number of processors must be less than or equal to the number of quadrature points.", CURRENT_FUNCTION ); } - if( _numProcs == 1 ) { - _localNSys = _nSys; - _startSysIdx = 0; - _endSysIdx = _nSys; - } - else { - _localNSys = _nSys / ( _numProcs - 1 ); - _startSysIdx = _rank * _localNSys; - _endSysIdx = _rank * _localNSys + _localNSys; - - if( _rank == _numProcs - 1 ) { - _localNSys = _nSys - _startSysIdx; - _endSysIdx = _nSys; - } - } + const unsigned long numRanks = static_cast( _numProcs ); + const unsigned long rankIndex = static_cast( _rank ); + const unsigned long baseChunk = _nSys / numRanks; + const unsigned long remainder = _nSys % numRanks; + + _localNSys = baseChunk + ( rankIndex < remainder ? 1UL : 0UL ); + _startSysIdx = rankIndex * baseChunk + std::min( rankIndex, remainder ); + _endSysIdx = _startSysIdx + _localNSys; // std::cout << "Rank: " << _rank << " startSysIdx: " << _startSysIdx << " endSysIdx: " << _endSysIdx << " localNSys: " << _localNSys << // std::endl; @@ -613,9 +606,29 @@ void SNSolverHPCCUDA::InitCUDA() { ErrorMessages::Error( "No CUDA-capable GPU detected, but SNSolverHPCCUDA was requested.", CURRENT_FUNCTION ); } - _cudaDeviceId = 0; // first version: pin to one GPU + int localRank = 0; + int localSize = 1; +#ifdef IMPORT_MPI + MPI_Comm localComm = MPI_COMM_NULL; + MPI_Comm_split_type( MPI_COMM_WORLD, MPI_COMM_TYPE_SHARED, _rank, MPI_INFO_NULL, &localComm ); + MPI_Comm_rank( localComm, &localRank ); + MPI_Comm_size( localComm, &localSize ); + MPI_Comm_free( &localComm ); +#endif + + _cudaDeviceId = localRank % nDevices; CheckCuda( cudaSetDevice( _cudaDeviceId ), "cudaSetDevice" ); + if( _rank == 0 ) { + auto log = spdlog::get( "event" ); + if( log ) { + log->info( "| CUDA backend: {} local MPI rank(s), {} visible CUDA device(s).", localSize, nDevices ); + if( localSize > nDevices ) { + log->warn( "| CUDA backend: {} local MPI rank(s) exceed {} visible device(s); GPUs will be shared.", localSize, nDevices ); + } + } + } + _device = new DeviceBuffers(); const std::size_t nCells = static_cast( _nCells ); @@ -805,6 +818,20 @@ void SNSolverHPCCUDA::Solve() { RK2AverageAndScalarFluxKernel<<>>( _nCells, _localNSys, _device->quadWeights, _device->solRK0, _device->sol, _device->scalarFlux ); CheckCuda( cudaGetLastError(), "RK2AverageAndScalarFluxKernel launch" ); +#ifdef IMPORT_MPI + CheckCuda( cudaMemcpy( _scalarFlux.data(), + _device->scalarFlux, + static_cast( _nCells ) * sizeof( double ), + cudaMemcpyDeviceToHost ), + "download scalar flux after RK2 average" ); + std::vector tempScalarFlux( _scalarFlux ); + MPI_Barrier( MPI_COMM_WORLD ); + MPI_Allreduce( tempScalarFlux.data(), _scalarFlux.data(), _nCells, MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD ); + MPI_Barrier( MPI_COMM_WORLD ); + CheckCuda( + cudaMemcpy( _device->scalarFlux, _scalarFlux.data(), static_cast( _nCells ) * sizeof( double ), cudaMemcpyHostToDevice ), + "sync allreduced scalar flux after RK2 average" ); +#endif } else { ( _spatialOrder == 2 ) ? FluxOrder2() : FluxOrder1(); diff --git a/tools/singularity/install_kitrt_singularity_cuda.sh b/tools/singularity/install_kitrt_singularity_cuda.sh index b88fe4a1..a59653d2 100755 --- a/tools/singularity/install_kitrt_singularity_cuda.sh +++ b/tools/singularity/install_kitrt_singularity_cuda.sh @@ -4,5 +4,5 @@ set -euo pipefail cd ../../ mkdir -p build_singularity_cuda cd build_singularity_cuda -cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_MPI=OFF -DBUILD_CUDA_HPC=ON -DBUILD_ML=OFF .. +cmake -DCMAKE_BUILD_TYPE=Release -DBUILD_MPI=ON -DBUILD_CUDA_HPC=ON -DBUILD_ML=OFF .. make -j From b308cf06331efd98b0808450a37c0582827fc778 Mon Sep 17 00:00:00 2001 From: Steffen Date: Wed, 18 Mar 2026 14:42:51 -0400 Subject: [PATCH 2/2] fix avg_block issue --- include/common/globalconstants.hpp | 4 +++ include/solvers/snsolver_hpc.hpp | 3 ++ include/solvers/snsolver_hpc_cuda.hpp | 3 ++ include/solvers/snsolver_hpc_hip.hpp | 3 ++ src/common/config.cpp | 13 ++++++-- src/solvers/snsolver_hpc.cpp | 43 +++++++++++++++++++++++++-- src/solvers/snsolver_hpc.cu | 43 +++++++++++++++++++++++++-- src/solvers/snsolver_hpc.hip | 43 +++++++++++++++++++++++++-- 8 files changed, 143 insertions(+), 12 deletions(-) diff --git a/include/common/globalconstants.hpp b/include/common/globalconstants.hpp index 4dca338f..bf50c726 100644 --- a/include/common/globalconstants.hpp +++ b/include/common/globalconstants.hpp @@ -199,6 +199,8 @@ enum SCALAR_OUTPUT { PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, ABSORPTION_GREEN_BLOCK, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_LINE }; @@ -225,6 +227,8 @@ inline std::map ScalarOutput_Map{ { "ITER", ITER }, { "PROBE_MOMENT_TIME_TRACE", PROBE_MOMENT_TIME_TRACE }, { "VAR_ABSORPTION_GREEN", VAR_ABSORPTION_GREEN }, { "ABSORPTION_GREEN_BLOCK", ABSORPTION_GREEN_BLOCK }, + { "AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED", AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED }, + { "VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED", VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED }, { "ABSORPTION_GREEN_LINE", ABSORPTION_GREEN_LINE } }; // Spherical Basis Name diff --git a/include/solvers/snsolver_hpc.hpp b/include/solvers/snsolver_hpc.hpp index c9753b41..d8816dc4 100644 --- a/include/solvers/snsolver_hpc.hpp +++ b/include/solvers/snsolver_hpc.hpp @@ -113,6 +113,8 @@ class SNSolverHPC double _curAbsorptionHohlraumVertical; double _curAbsorptionHohlraumHorizontal; double _varAbsorptionHohlraumGreen; + double _avgAbsorptionHohlraumGreenBlockIntegrated; + double _varAbsorptionHohlraumGreenBlockIntegrated; std::vector> _probingCellsHohlraum; /*!< @brief Indices of cells that contain a probing sensor */ std::vector _probingMoments; /*!< @brief Solution Momnets at the probing cells that contain a probing sensor */ @@ -125,6 +127,7 @@ class SNSolverHPC unsigned _nProbingCellsBlocksGreen; std::vector> _probingCellsBlocksGreen; /*!< @brief Indices of cells that contain a probing sensor blocks */ std::vector _absorptionValsBlocksGreen; /*!< @brief Avg Absorption value at the sampleing blocks of lineGreen */ + std::vector _absorptionValsBlocksGreenIntegrated; /*!< @brief Internal QoI-3 running values A_i(t) = |g_i|^{-1} \int_{g_i}\int_0^t sigma_a phi d\tau dxdy */ // Design parameters std::vector _cornerUpperLeftGreen; /*!< @brief Coord of corner of the green area (minus thickness/2 of it) relative to the green center */ diff --git a/include/solvers/snsolver_hpc_cuda.hpp b/include/solvers/snsolver_hpc_cuda.hpp index 651d729e..c2ecc476 100644 --- a/include/solvers/snsolver_hpc_cuda.hpp +++ b/include/solvers/snsolver_hpc_cuda.hpp @@ -115,6 +115,8 @@ class SNSolverHPCCUDA double _curAbsorptionHohlraumVertical; double _curAbsorptionHohlraumHorizontal; double _varAbsorptionHohlraumGreen; + double _avgAbsorptionHohlraumGreenBlockIntegrated; + double _varAbsorptionHohlraumGreenBlockIntegrated; std::vector> _probingCellsHohlraum; /*!< @brief Indices of cells that contain a probing sensor */ std::vector _probingMoments; /*!< @brief Solution Momnets at the probing cells that contain a probing sensor */ @@ -127,6 +129,7 @@ class SNSolverHPCCUDA unsigned _nProbingCellsBlocksGreen; std::vector> _probingCellsBlocksGreen; /*!< @brief Indices of cells that contain a probing sensor blocks */ std::vector _absorptionValsBlocksGreen; /*!< @brief Avg Absorption value at the sampleing blocks of lineGreen */ + std::vector _absorptionValsBlocksGreenIntegrated; /*!< @brief Internal QoI-3 running values A_i(t) = |g_i|^{-1} \int_{g_i}\int_0^t sigma_a phi d\tau dxdy */ // Design parameters std::vector _cornerUpperLeftGreen; /*!< @brief Coord of corner of the green area (minus thickness/2 of it) relative to the green center */ diff --git a/include/solvers/snsolver_hpc_hip.hpp b/include/solvers/snsolver_hpc_hip.hpp index 6789b423..7a6bc7af 100644 --- a/include/solvers/snsolver_hpc_hip.hpp +++ b/include/solvers/snsolver_hpc_hip.hpp @@ -115,6 +115,8 @@ class SNSolverHPCHIP double _curAbsorptionHohlraumVertical; double _curAbsorptionHohlraumHorizontal; double _varAbsorptionHohlraumGreen; + double _avgAbsorptionHohlraumGreenBlockIntegrated; + double _varAbsorptionHohlraumGreenBlockIntegrated; std::vector> _probingCellsHohlraum; /*!< @brief Indices of cells that contain a probing sensor */ std::vector _probingMoments; /*!< @brief Solution Momnets at the probing cells that contain a probing sensor */ @@ -127,6 +129,7 @@ class SNSolverHPCHIP unsigned _nProbingCellsBlocksGreen; std::vector> _probingCellsBlocksGreen; /*!< @brief Indices of cells that contain a probing sensor blocks */ std::vector _absorptionValsBlocksGreen; /*!< @brief Avg Absorption value at the sampleing blocks of lineGreen */ + std::vector _absorptionValsBlocksGreenIntegrated; /*!< @brief Internal QoI-3 running values A_i(t) = |g_i|^{-1} \int_{g_i}\int_0^t sigma_a phi d\tau dxdy */ // Design parameters std::vector _cornerUpperLeftGreen; /*!< @brief Coord of corner of the green area (minus thickness/2 of it) relative to the green center */ diff --git a/src/common/config.cpp b/src/common/config.cpp index 246efdb3..7b50d828 100644 --- a/src/common/config.cpp +++ b/src/common/config.cpp @@ -806,6 +806,8 @@ void Config::SetPostprocessing() { TOTAL_PARTICLE_ABSORPTION, PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, }; it = std::find( legalOutputs.begin(), legalOutputs.end(), _screenOutput[idx_screenOutput] ); @@ -813,11 +815,12 @@ void Config::SetPostprocessing() { if( it == legalOutputs.end() ) { std::string foundKey = findKey( ScalarOutput_Map, _screenOutput[idx_screenOutput] ); ErrorMessages::Error( - "Illegal output field <" + foundKey + + "Illegal output field <" + foundKey + "> for option SCREEN_OUTPUT for this test case.\n" "Supported fields are: ITER, SIM_TIME, WALL_TIME, MASS, RMS_FLUX, VTK_OUTPUT, CSV_OUTPUT, TOTAL_PARTICLE_ABSORPTION_CENTER, \n" "TOTAL_PARTICLE_ABSORPTION_VERTICAL, TOTAL_PARTICLE_ABSORPTION_HORIZONTAL, PROBE_MOMENT_TIME_TRACE, CUR_OUTFLOW, \n " - "TOTAL_OUTFLOW, MAX_OUTFLOW, VAR_ABSORPTION_GREEN \n" + "TOTAL_OUTFLOW, MAX_OUTFLOW, VAR_ABSORPTION_GREEN, AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, " + "VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED \n" "Please check your .cfg file.", CURRENT_FUNCTION ); } @@ -957,6 +960,8 @@ void Config::SetPostprocessing() { PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, ABSORPTION_GREEN_BLOCK, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_LINE }; it = std::find( legalOutputs.begin(), legalOutputs.end(), _historyOutput[idx_historyOutput] ); @@ -968,7 +973,8 @@ void Config::SetPostprocessing() { "> for option HISTORY_OUTPUT for this test case.\n" "Supported fields are: ITER, SIM_TIME, WALL_TIME, MASS RMS_FLUX, VTK_OUTPUT, CSV_OUTPUT, TOTAL_PARTICLE_ABSORPTION_CENTER, \n " "TOTAL_PARTICLE_ABSORPTION_VERTICAL, TOTAL_PARTICLE_ABSORPTION_HORIZONTAL,PROBE_MOMENT_TIME_TRACE, CUR_OUTFLOW, \n" - "TOTAL_OUTFLOW, MAX_OUTFLOW , VAR_ABSORPTION_GREEN, ABSORPTION_GREEN_BLOCK, ABSORPTION_GREEN_LINE \n" + "TOTAL_OUTFLOW, MAX_OUTFLOW , VAR_ABSORPTION_GREEN, ABSORPTION_GREEN_BLOCK, " + "AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_LINE \n" "Please check your .cfg file.", CURRENT_FUNCTION ); } @@ -1043,6 +1049,7 @@ void Config::SetPostprocessing() { _nHistoryOutput += 44 - 1; // extend the screen output by the number of probing points for( unsigned i = 0; i < 44; i++ ) _historyOutput.push_back( ABSORPTION_GREEN_BLOCK ); } + } } diff --git a/src/solvers/snsolver_hpc.cpp b/src/solvers/snsolver_hpc.cpp index 678a78d3..dd284d6f 100644 --- a/src/solvers/snsolver_hpc.cpp +++ b/src/solvers/snsolver_hpc.cpp @@ -193,6 +193,8 @@ SNSolverHPC::SNSolverHPC( Config* settings ) { _curAbsorptionHohlraumVertical = 0; _curAbsorptionHohlraumHorizontal = 0; _varAbsorptionHohlraumGreen = 0; + _avgAbsorptionHohlraumGreenBlockIntegrated = 0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0; } if( _settings->GetLoadRestartSolution() ) @@ -267,9 +269,10 @@ SNSolverHPC::SNSolverHPC( Config* settings ) { _nProbingCellsLineGreen = _settings->GetNumProbingCellsLineHohlraum(); - _nProbingCellsBlocksGreen = 44; - _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); - _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); + _nProbingCellsBlocksGreen = 44; + _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsBlocksGreenIntegrated = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); SetProbingCellsLineGreen(); // ONLY FOR HOHLRAUM } @@ -825,6 +828,8 @@ void SNSolverHPC::IterPostprocessing() { } // Update time integral values on rank 0 if( _rank == 0 ) { + constexpr double greenBlockArea = 0.05 * 0.05; + _totalScalarOutflow += _curScalarOutflow * _dT; _totalScalarOutflowPeri1 += _curScalarOutflowPeri1 * _dT; _totalScalarOutflowPeri2 += _curScalarOutflowPeri2 * _dT; @@ -833,6 +838,20 @@ void SNSolverHPC::IterPostprocessing() { _totalAbsorptionHohlraumCenter += _curAbsorptionHohlraumCenter * _dT; _totalAbsorptionHohlraumVertical += _curAbsorptionHohlraumVertical * _dT; _totalAbsorptionHohlraumHorizontal += _curAbsorptionHohlraumHorizontal * _dT; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _absorptionValsBlocksGreenIntegrated[i] += _dT * _absorptionValsBlocksGreen[i] / greenBlockArea; + } + _avgAbsorptionHohlraumGreenBlockIntegrated = 0.0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0.0; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _avgAbsorptionHohlraumGreenBlockIntegrated += _absorptionValsBlocksGreenIntegrated[i]; + } + _avgAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + const double diff = _absorptionValsBlocksGreenIntegrated[i] - _avgAbsorptionHohlraumGreenBlockIntegrated; + _varAbsorptionHohlraumGreenBlockIntegrated += diff * diff; + } + _varAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); _rmsFlux = sqrt( _rmsFlux ); } @@ -938,6 +957,8 @@ void SNSolverHPC::PrepareScreenOutput() { } break; case VAR_ABSORPTION_GREEN: _screenOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "V_G"; break; default: ErrorMessages::Error( "Screen output field not defined!", CURRENT_FUNCTION ); break; } @@ -996,6 +1017,12 @@ void SNSolverHPC::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; default: ErrorMessages::Error( "Screen output group not defined!", CURRENT_FUNCTION ); break; } } @@ -1054,6 +1081,12 @@ void SNSolverHPC::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; case ABSORPTION_GREEN_LINE: for( unsigned i = 0; i < _settings->GetNumProbingCellsLineHohlraum(); i++ ) { _historyOutputFields[idx_field] = _absorptionValsLineSegment[i]; @@ -1106,6 +1139,8 @@ void SNSolverHPC::PrintScreenOutput( unsigned idx_iter ) { TOTAL_PARTICLE_ABSORPTION_HORIZONTAL, PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_BLOCK, ABSORPTION_GREEN_LINE }; std::vector booleanFields = { VTK_OUTPUT, CSV_OUTPUT }; @@ -1187,6 +1222,8 @@ void SNSolverHPC::PrepareHistoryOutput() { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "V_G"; break; case ABSORPTION_GREEN_BLOCK: for( unsigned i = 0; i < 44; i++ ) { _historyOutputFieldNames[idx_field] = "Probe Green Block " + std::to_string( i ); diff --git a/src/solvers/snsolver_hpc.cu b/src/solvers/snsolver_hpc.cu index c6facfca..c9f5ddfd 100644 --- a/src/solvers/snsolver_hpc.cu +++ b/src/solvers/snsolver_hpc.cu @@ -513,6 +513,8 @@ SNSolverHPCCUDA::SNSolverHPCCUDA( Config* settings ) { _curAbsorptionHohlraumVertical = 0; _curAbsorptionHohlraumHorizontal = 0; _varAbsorptionHohlraumGreen = 0; + _avgAbsorptionHohlraumGreenBlockIntegrated = 0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0; } if( _settings->GetLoadRestartSolution() ) @@ -591,9 +593,10 @@ SNSolverHPCCUDA::SNSolverHPCCUDA( Config* settings ) { _nProbingCellsLineGreen = _settings->GetNumProbingCellsLineHohlraum(); - _nProbingCellsBlocksGreen = 44; - _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); - _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); + _nProbingCellsBlocksGreen = 44; + _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsBlocksGreenIntegrated = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); SetProbingCellsLineGreen(); // ONLY FOR HOHLRAUM } @@ -1223,6 +1226,8 @@ void SNSolverHPCCUDA::IterPostprocessing() { } // Update time integral values on rank 0 if( _rank == 0 ) { + constexpr double greenBlockArea = 0.05 * 0.05; + _totalScalarOutflow += _curScalarOutflow * _dT; _totalScalarOutflowPeri1 += _curScalarOutflowPeri1 * _dT; _totalScalarOutflowPeri2 += _curScalarOutflowPeri2 * _dT; @@ -1231,6 +1236,20 @@ void SNSolverHPCCUDA::IterPostprocessing() { _totalAbsorptionHohlraumCenter += _curAbsorptionHohlraumCenter * _dT; _totalAbsorptionHohlraumVertical += _curAbsorptionHohlraumVertical * _dT; _totalAbsorptionHohlraumHorizontal += _curAbsorptionHohlraumHorizontal * _dT; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _absorptionValsBlocksGreenIntegrated[i] += _dT * _absorptionValsBlocksGreen[i] / greenBlockArea; + } + _avgAbsorptionHohlraumGreenBlockIntegrated = 0.0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0.0; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _avgAbsorptionHohlraumGreenBlockIntegrated += _absorptionValsBlocksGreenIntegrated[i]; + } + _avgAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + const double diff = _absorptionValsBlocksGreenIntegrated[i] - _avgAbsorptionHohlraumGreenBlockIntegrated; + _varAbsorptionHohlraumGreenBlockIntegrated += diff * diff; + } + _varAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); _rmsFlux = sqrt( _rmsFlux ); } @@ -1336,6 +1355,8 @@ void SNSolverHPCCUDA::PrepareScreenOutput() { } break; case VAR_ABSORPTION_GREEN: _screenOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "V_G"; break; default: ErrorMessages::Error( "Screen output field not defined!", CURRENT_FUNCTION ); break; } @@ -1394,6 +1415,12 @@ void SNSolverHPCCUDA::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; default: ErrorMessages::Error( "Screen output group not defined!", CURRENT_FUNCTION ); break; } } @@ -1452,6 +1479,12 @@ void SNSolverHPCCUDA::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; case ABSORPTION_GREEN_LINE: for( unsigned i = 0; i < _settings->GetNumProbingCellsLineHohlraum(); i++ ) { _historyOutputFields[idx_field] = _absorptionValsLineSegment[i]; @@ -1504,6 +1537,8 @@ void SNSolverHPCCUDA::PrintScreenOutput( unsigned idx_iter ) { TOTAL_PARTICLE_ABSORPTION_HORIZONTAL, PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_BLOCK, ABSORPTION_GREEN_LINE }; std::vector booleanFields = { VTK_OUTPUT, CSV_OUTPUT }; @@ -1585,6 +1620,8 @@ void SNSolverHPCCUDA::PrepareHistoryOutput() { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "V_G"; break; case ABSORPTION_GREEN_BLOCK: for( unsigned i = 0; i < 44; i++ ) { _historyOutputFieldNames[idx_field] = "Probe Green Block " + std::to_string( i ); diff --git a/src/solvers/snsolver_hpc.hip b/src/solvers/snsolver_hpc.hip index 3418da0b..990134e0 100644 --- a/src/solvers/snsolver_hpc.hip +++ b/src/solvers/snsolver_hpc.hip @@ -507,6 +507,8 @@ SNSolverHPCHIP::SNSolverHPCHIP( Config* settings ) { _curAbsorptionHohlraumVertical = 0; _curAbsorptionHohlraumHorizontal = 0; _varAbsorptionHohlraumGreen = 0; + _avgAbsorptionHohlraumGreenBlockIntegrated = 0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0; } if( _settings->GetLoadRestartSolution() ) @@ -585,9 +587,10 @@ SNSolverHPCHIP::SNSolverHPCHIP( Config* settings ) { _nProbingCellsLineGreen = _settings->GetNumProbingCellsLineHohlraum(); - _nProbingCellsBlocksGreen = 44; - _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); - _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); + _nProbingCellsBlocksGreen = 44; + _absorptionValsBlocksGreen = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsBlocksGreenIntegrated = std::vector( _nProbingCellsBlocksGreen, 0. ); + _absorptionValsLineSegment = std::vector( _nProbingCellsLineGreen, 0.0 ); SetProbingCellsLineGreen(); // ONLY FOR HOHLRAUM } @@ -1213,6 +1216,8 @@ void SNSolverHPCHIP::IterPostprocessing() { } // Update time integral values on rank 0 if( _rank == 0 ) { + constexpr double greenBlockArea = 0.05 * 0.05; + _totalScalarOutflow += _curScalarOutflow * _dT; _totalScalarOutflowPeri1 += _curScalarOutflowPeri1 * _dT; _totalScalarOutflowPeri2 += _curScalarOutflowPeri2 * _dT; @@ -1221,6 +1226,20 @@ void SNSolverHPCHIP::IterPostprocessing() { _totalAbsorptionHohlraumCenter += _curAbsorptionHohlraumCenter * _dT; _totalAbsorptionHohlraumVertical += _curAbsorptionHohlraumVertical * _dT; _totalAbsorptionHohlraumHorizontal += _curAbsorptionHohlraumHorizontal * _dT; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _absorptionValsBlocksGreenIntegrated[i] += _dT * _absorptionValsBlocksGreen[i] / greenBlockArea; + } + _avgAbsorptionHohlraumGreenBlockIntegrated = 0.0; + _varAbsorptionHohlraumGreenBlockIntegrated = 0.0; + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + _avgAbsorptionHohlraumGreenBlockIntegrated += _absorptionValsBlocksGreenIntegrated[i]; + } + _avgAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); + for( unsigned i = 0; i < _nProbingCellsBlocksGreen; ++i ) { + const double diff = _absorptionValsBlocksGreenIntegrated[i] - _avgAbsorptionHohlraumGreenBlockIntegrated; + _varAbsorptionHohlraumGreenBlockIntegrated += diff * diff; + } + _varAbsorptionHohlraumGreenBlockIntegrated /= static_cast( _nProbingCellsBlocksGreen ); _rmsFlux = sqrt( _rmsFlux ); } @@ -1326,6 +1345,8 @@ void SNSolverHPCHIP::PrepareScreenOutput() { } break; case VAR_ABSORPTION_GREEN: _screenOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _screenOutputFieldNames[idx_field] = "V_G"; break; default: ErrorMessages::Error( "Screen output field not defined!", CURRENT_FUNCTION ); break; } @@ -1384,6 +1405,12 @@ void SNSolverHPCHIP::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _screenOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; default: ErrorMessages::Error( "Screen output group not defined!", CURRENT_FUNCTION ); break; } } @@ -1442,6 +1469,12 @@ void SNSolverHPCHIP::WriteScalarOutput( unsigned idx_iter ) { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreen; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _avgAbsorptionHohlraumGreenBlockIntegrated; + break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: + _historyOutputFields[idx_field] = _varAbsorptionHohlraumGreenBlockIntegrated; + break; case ABSORPTION_GREEN_LINE: for( unsigned i = 0; i < _settings->GetNumProbingCellsLineHohlraum(); i++ ) { _historyOutputFields[idx_field] = _absorptionValsLineSegment[i]; @@ -1494,6 +1527,8 @@ void SNSolverHPCHIP::PrintScreenOutput( unsigned idx_iter ) { TOTAL_PARTICLE_ABSORPTION_HORIZONTAL, PROBE_MOMENT_TIME_TRACE, VAR_ABSORPTION_GREEN, + AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED, + VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED, ABSORPTION_GREEN_BLOCK, ABSORPTION_GREEN_LINE }; std::vector booleanFields = { VTK_OUTPUT, CSV_OUTPUT }; @@ -1575,6 +1610,8 @@ void SNSolverHPCHIP::PrepareHistoryOutput() { idx_field--; break; case VAR_ABSORPTION_GREEN: _historyOutputFieldNames[idx_field] = "Var. absorption green"; break; + case AVG_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "A_G"; break; + case VAR_ABSORPTION_GREEN_BLOCK_INTEGRATED: _historyOutputFieldNames[idx_field] = "V_G"; break; case ABSORPTION_GREEN_BLOCK: for( unsigned i = 0; i < 44; i++ ) { _historyOutputFieldNames[idx_field] = "Probe Green Block " + std::to_string( i );