diff --git a/src/common/pmpl.hpp b/src/common/pmpl.hpp index ce622cd..b8a0fbe 100644 --- a/src/common/pmpl.hpp +++ b/src/common/pmpl.hpp @@ -22,17 +22,17 @@ namespace hpcReact { #if defined(HPCREACT_USE_DEVICE) #if defined(HPCREACT_USE_CUDA) - #define deviceMalloc( PTR, BYTES ) cudaMalloc( PTR, BYTES ); - #define deviceMallocManaged( PTR, BYTES ) cudaMallocManaged( PTR, BYTES ); - #define deviceDeviceSynchronize() cudaDeviceSynchronize(); - #define deviceMemCpy( DST, SRC, BYTES, KIND ) cudaMemcpy( DST, SRC, BYTES, KIND ); - #define deviceFree( PTR ) cudaFree( PTR ); +#define deviceMalloc( PTR, BYTES ) cudaMalloc( PTR, BYTES ); +#define deviceMallocManaged( PTR, BYTES ) cudaMallocManaged( PTR, BYTES ); +#define deviceDeviceSynchronize() cudaDeviceSynchronize(); +#define deviceMemCpy( DST, SRC, BYTES, KIND ) cudaMemcpy( DST, SRC, BYTES, KIND ); +#define deviceFree( PTR ) cudaFree( PTR ); #elif defined(HPCREACT_USE_HIP) - #define deviceMalloc( PTR, BYTES ) hipMalloc( PTR, BYTES ); - #define deviceMallocManaged( PTR, BYTES ) hipMallocManaged( PTR, BYTES ); - #define deviceDeviceSynchronize() hipDeviceSynchronize(); - #define deviceMemCpy( DST, SRC, BYTES, KIND ) hipMemcpy( DST, SRC, BYTES, KIND ); - #define deviceFree( PTR ) hipFree( PTR ); +#define deviceMalloc( PTR, BYTES ) hipMalloc( PTR, BYTES ); +#define deviceMallocManaged( PTR, BYTES ) hipMallocManaged( PTR, BYTES ); +#define deviceDeviceSynchronize() hipDeviceSynchronize(); +#define deviceMemCpy( DST, SRC, BYTES, KIND ) hipMemcpy( DST, SRC, BYTES, KIND ); +#define deviceFree( PTR ) hipFree( PTR ); #endif #endif @@ -130,12 +130,18 @@ void genericKernelWrapper( int const N, DATA_TYPE * const hostData, LAMBDA && fu genericKernel <<< 1, 1 >>> ( std::forward< LAMBDA >( func ), deviceData ); cudaError_t e = cudaGetLastError(); - if (e != cudaSuccess) { fprintf(stderr, "launch error: %s\n", cudaGetErrorString(e)); abort(); } + if( e != cudaSuccess ) + { + fprintf( stderr, "launch error: %s\n", cudaGetErrorString( e )); abort(); + } deviceDeviceSynchronize(); e = cudaGetLastError(); - if (e != cudaSuccess) { fprintf(stderr, "post-sync error: %s\n", cudaGetErrorString(e)); abort(); } + if( e != cudaSuccess ) + { + fprintf( stderr, "post-sync error: %s\n", cudaGetErrorString( e )); abort(); + } deviceMemCpy( hostData, deviceData, N * sizeof(DATA_TYPE), cudaMemcpyDeviceToHost ); deviceFree( deviceData ); diff --git a/src/reactions/exampleSystems/unitTests/testMomasEasyCase.cpp b/src/reactions/exampleSystems/unitTests/testMomasEasyCase.cpp index 33b4669..23ad284 100644 --- a/src/reactions/exampleSystems/unitTests/testMomasEasyCase.cpp +++ b/src/reactions/exampleSystems/unitTests/testMomasEasyCase.cpp @@ -65,7 +65,7 @@ void testMoMasAllEquilibriumHelper() targetAggregatePrimarySpeciesConcentration, logInitialPrimarySpeciesConcentration, logPrimarySpeciesConcentrationCopy ); - }); + } ); double const expectedPrimarySpeciesConcentrations[numPrimarySpecies] = { diff --git a/src/reactions/exampleSystems/unitTests/testMomasMediumCase.cpp b/src/reactions/exampleSystems/unitTests/testMomasMediumCase.cpp index 534cd03..880f9b9 100644 --- a/src/reactions/exampleSystems/unitTests/testMomasMediumCase.cpp +++ b/src/reactions/exampleSystems/unitTests/testMomasMediumCase.cpp @@ -60,13 +60,13 @@ void testMoMasMediumEquilibriumHelper() log( initialPrimarySpeciesConcentration[3] ), log( initialPrimarySpeciesConcentration[4] ) }; - + EquilibriumReactionsType::enforceEquilibrium_Aggregate( 0, hpcReact::MoMasBenchmark::mediumCaseParams.equilibriumReactionsParameters(), targetAggregatePrimarySpeciesConcentration, logInitialPrimarySpeciesConcentration, logPrimarySpeciesConcentrationCopy ); - }); + } ); double const expectedPrimarySpeciesConcentrations[numPrimarySpecies] = { diff --git a/src/reactions/massActions/MassActions.hpp b/src/reactions/massActions/MassActions.hpp index 17dbbf8..41b8183 100644 --- a/src/reactions/massActions/MassActions.hpp +++ b/src/reactions/massActions/MassActions.hpp @@ -82,7 +82,7 @@ void calculateLogSecondarySpeciesConcentration( PARAMS_DATA const & params, INDEX_TYPE >( params, logPrimarySpeciesConcentrations, logSecondarySpeciesConcentrations, - [](INDEX_TYPE, INDEX_TYPE, REAL_TYPE ){} ); + []( INDEX_TYPE, INDEX_TYPE, REAL_TYPE ){} ); } @@ -154,8 +154,7 @@ void calculateAggregatePrimaryConcentrationsWrtLogC( PARAMS_DATA const & params, for( int k=0; k data; + CalculateLogSecondarySpeciesConcentrationData< numPrimarySpecies, numSecondarySpecies > data; pmpl::genericKernelWrapper( 1, &data, [] HPCREACT_DEVICE ( auto * const dataCopy ) { calculateLogSecondarySpeciesConcentration< double, - int, - int >( carbonateSystemAllEquilibrium.equilibriumReactionsParameters(), + int, + int >( carbonateSystemAllEquilibrium.equilibriumReactionsParameters(), dataCopy->logPrimarySpeciesSolution, dataCopy->logSecondarySpeciesConcentrations ); @@ -60,11 +60,10 @@ void test_calculateLogSecondarySpeciesConcentration_helper() calculateLogSecondarySpeciesConcentrationWrtLogC< double, int, int >( carbonateSystemAllEquilibrium.equilibriumReactionsParameters(), - dataCopy->logPrimarySpeciesSolution, - dataCopy->logSecondarySpeciesConcentrations, - dataCopy->dLogSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentrations ); - }); - + dataCopy->logPrimarySpeciesSolution, + dataCopy->logSecondarySpeciesConcentrations, + dataCopy->dLogSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentrations ); + } ); @@ -126,13 +125,13 @@ struct CalculateAggregatePrimaryConcentrationsWrtLogCHelperData { double const primarySpeciesSolution[numPrimarySpecies] = { - log(0.00043969547214915125), - log(0.00037230096984514874), - log(0.014716565308128551), - log(0.0024913722747387217), - log(1.8586090945989489), - log(0.009881874292035079), - log(1.0723078278653704) + log( 0.00043969547214915125 ), + log( 0.00037230096984514874 ), + log( 0.014716565308128551 ), + log( 0.0024913722747387217 ), + log( 1.8586090945989489 ), + log( 0.009881874292035079 ), + log( 1.0723078278653704 ) }; double aggregatePrimarySpeciesConcentration[numPrimarySpecies] = {0}; @@ -144,7 +143,7 @@ void testcalculateAggregatePrimaryConcentrationsWrtLogCHelper() { static constexpr int numPrimarySpecies = carbonateSystemAllEquilibrium.numPrimarySpecies(); - CalculateAggregatePrimaryConcentrationsWrtLogCHelperData data; + CalculateAggregatePrimaryConcentrationsWrtLogCHelperData< numPrimarySpecies > data; pmpl::genericKernelWrapper( 1, &data, [] HPCREACT_DEVICE ( auto * const dataCopy ) { @@ -152,7 +151,7 @@ void testcalculateAggregatePrimaryConcentrationsWrtLogCHelper() dataCopy->primarySpeciesSolution, dataCopy->aggregatePrimarySpeciesConcentration, dataCopy->dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations ); - }); + } ); double const expectedAggregatePrimarySpeciesConcentration[numPrimarySpecies] = { diff --git a/src/reactions/reactionsSystems/Parameters.hpp b/src/reactions/reactionsSystems/Parameters.hpp index 54cddaa..d28d37a 100644 --- a/src/reactions/reactionsSystems/Parameters.hpp +++ b/src/reactions/reactionsSystems/Parameters.hpp @@ -157,7 +157,7 @@ struct MixedReactionsParameters HPCREACT_HOST_DEVICE static constexpr IndexType numSecondarySpecies() { return NUM_EQ_REACTIONS; } - HPCREACT_HOST_DEVICE + HPCREACT_HOST_DEVICE constexpr EquilibriumReactionsParameters< RealType, IntType, IndexType, numSpecies(), numEquilibriumReactions() > equilibriumReactionsParameters() const @@ -179,7 +179,7 @@ struct MixedReactionsParameters return { eqMatrix, eqConstants, mobileSpeciesFlags }; } - HPCREACT_HOST_DEVICE + HPCREACT_HOST_DEVICE constexpr KineticReactionsParameters< RealType, IntType, IndexType, numSpecies(), numKineticReactions() > kineticReactionsParameters() const @@ -203,7 +203,7 @@ struct MixedReactionsParameters return { kineticMatrix, rateConstantForward, rateConstantReverse, equilibriumConstant, m_reactionRatesUpdateOption }; } - HPCREACT_HOST_DEVICE + HPCREACT_HOST_DEVICE void verifyParameterConsistency() { static constexpr int num_digits = 12; @@ -239,10 +239,10 @@ struct MixedReactionsParameters } } - HPCREACT_HOST_DEVICE RealType stoichiometricMatrix( IndexType const r, int const i ) const { return m_stoichiometricMatrix[r][i]; } - HPCREACT_HOST_DEVICE RealType equilibriumConstant( IndexType const r ) const { return m_equilibriumConstant[r]; } - HPCREACT_HOST_DEVICE RealType rateConstantForward( IndexType const r ) const { return m_rateConstantForward[r]; } - HPCREACT_HOST_DEVICE RealType rateConstantReverse( IndexType const r ) const { return m_rateConstantReverse[r]; } + HPCREACT_HOST_DEVICE RealType stoichiometricMatrix( IndexType const r, int const i ) const { return m_stoichiometricMatrix[r][i]; } + HPCREACT_HOST_DEVICE RealType equilibriumConstant( IndexType const r ) const { return m_equilibriumConstant[r]; } + HPCREACT_HOST_DEVICE RealType rateConstantForward( IndexType const r ) const { return m_rateConstantForward[r]; } + HPCREACT_HOST_DEVICE RealType rateConstantReverse( IndexType const r ) const { return m_rateConstantReverse[r]; } CArrayWrapper< RealType, NUM_REACTIONS, NUM_SPECIES > m_stoichiometricMatrix; CArrayWrapper< RealType, NUM_REACTIONS > m_equilibriumConstant; diff --git a/src/reactions/unitTestUtilities/equilibriumReactionsTestUtilities.hpp b/src/reactions/unitTestUtilities/equilibriumReactionsTestUtilities.hpp index 58b0bdb..566eca3 100644 --- a/src/reactions/unitTestUtilities/equilibriumReactionsTestUtilities.hpp +++ b/src/reactions/unitTestUtilities/equilibriumReactionsTestUtilities.hpp @@ -35,7 +35,7 @@ REAL_TYPE tolerance( REAL_TYPE const a, REAL_TYPE const b ) * @tparam numSpecies Number of species. */ template< int numReactions, int numSpecies > -struct ComputeResidualAndJacobianTestData +struct ComputeResidualAndJacobianTestData { /// The reaction residuals double residual[numReactions] = { 0.0 }; @@ -65,23 +65,23 @@ void computeResidualAndJacobianTest( PARAMS_DATA const & params, double const temperature = 298.15; - ComputeResidualAndJacobianTestData data; + ComputeResidualAndJacobianTestData< numReactions, numSpecies > data; for( int i = 0; i < numSpecies; ++i ) { data.speciesConcentration[i] = initialSpeciesConcentration[i]; - } + } pmpl::genericKernelWrapper( 1, &data, [params, temperature] HPCREACT_DEVICE ( auto * const dataCopy ) - { - double xi[numReactions] = { 0.0 }; + { + double xi[numReactions] = { 0.0 }; - EquilibriumReactionsType::computeResidualAndJacobianReactionExtents( temperature, - params, - dataCopy->speciesConcentration, - xi, - dataCopy->residual, - dataCopy->jacobian ); - }); + EquilibriumReactionsType::computeResidualAndJacobianReactionExtents( temperature, + params, + dataCopy->speciesConcentration, + xi, + dataCopy->residual, + dataCopy->jacobian ); + } ); // printf( "R = { %8.4g, %8.4g }\n", residual[0], residual[1] ); for( int r=0; r data; + TestEnforceEquilibriumData< numSpecies > data; for( int i = 0; i < numSpecies; ++i ) { data.speciesConcentration0[i] = initialSpeciesConcentration[i]; } pmpl::genericKernelWrapper( 1, &data, [params, temperature] HPCREACT_DEVICE ( auto * const dataCopy ) - { - EquilibriumReactionsType::enforceEquilibrium_Extents( temperature, - params, - dataCopy->speciesConcentration0, - dataCopy->speciesConcentration ); - }); + { + EquilibriumReactionsType::enforceEquilibrium_Extents( temperature, + params, + dataCopy->speciesConcentration0, + dataCopy->speciesConcentration ); + } ); for( int r=0; rspeciesConcentration, - dataCopy->surfaceArea, - dataCopy->reactionRates, - dataCopy->reactionRatesDerivatives ); - }); + { + KineticReactionsType::computeReactionRates( temperature, + params, + dataCopy->speciesConcentration, + dataCopy->surfaceArea, + dataCopy->reactionRates, + dataCopy->reactionRatesDerivatives ); + } ); for( int r=0; rspeciesConcentration, - dataCopy->speciesRates, - dataCopy->speciesRatesDerivatives ); - }); + { + KineticReactionsType::computeSpeciesRates( temperature, + params, + dataCopy->speciesConcentration, + dataCopy->speciesRates, + dataCopy->speciesRatesDerivatives ); + } ); for( int r=0; r speciesRatesDerivatives; - - for( int t = 0; t < numSteps; ++t ) - { - printf("Time step %d \n ", t); - for( int i=0; ispeciesConcentration[i]; - } - KineticReactionsType::timeStep( dt, - temperature, - params, - speciesConcentration_n, - dataCopy->speciesConcentration, - speciesRates, - speciesRatesDerivatives ); - dataCopy->time += dt; - } - }); + double speciesConcentration_n[numSpecies]; + double speciesRates[numSpecies] = { 0.0 }; + CArrayWrapper< double, numSpecies, numSpecies > speciesRatesDerivatives; + + for( int t = 0; t < numSteps; ++t ) + { + printf( "Time step %d \n ", t ); + for( int i=0; ispeciesConcentration[i]; + } + KineticReactionsType::timeStep( dt, + temperature, + params, + speciesConcentration_n, + dataCopy->speciesConcentration, + speciesRates, + speciesRatesDerivatives ); + dataCopy->time += dt; + } + } ); EXPECT_NEAR( data.time, dt*numSteps, 1.0e-8 ); diff --git a/src/reactions/unitTestUtilities/mixedReactionsTestUtilities.hpp b/src/reactions/unitTestUtilities/mixedReactionsTestUtilities.hpp index 7426c03..2e39089 100644 --- a/src/reactions/unitTestUtilities/mixedReactionsTestUtilities.hpp +++ b/src/reactions/unitTestUtilities/mixedReactionsTestUtilities.hpp @@ -102,33 +102,33 @@ void timeStepTest( PARAMS_DATA const & params, } auto computeResidualAndJacobian = [&] ( REAL_TYPE const (&X)[numPrimarySpecies], - REAL_TYPE ( &r )[numPrimarySpecies], - REAL_TYPE ( &J )[numPrimarySpecies][numPrimarySpecies] ) + REAL_TYPE ( & r )[numPrimarySpecies], + REAL_TYPE ( & J )[numPrimarySpecies][numPrimarySpecies] ) + { + MixedReactionsType::updateMixedSystem( temperature, + params, + X, + surfaceArea, + logSecondarySpeciesConcentration, + aggregatePrimarySpeciesConcentration, + mobileAggregatePrimarySpeciesConcentration, + dAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration, + dMobileAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration, + reactionRates, + dReactionRates_dlogPrimarySpeciesConcentration, + aggregateSpeciesRates, + dAggregateSpeciesRates_dlogPrimarySpeciesConcentration ); + + + for( int i = 0; i < numPrimarySpecies; ++i ) + { + r[i] = ( aggregatePrimarySpeciesConcentration[i] - aggregatePrimarySpeciesConcentration_n[i] ) - aggregateSpeciesRates[i] * dt; + for( int j = 0; j < numPrimarySpecies; ++j ) { - MixedReactionsType::updateMixedSystem( temperature, - params, - X, - surfaceArea, - logSecondarySpeciesConcentration, - aggregatePrimarySpeciesConcentration, - mobileAggregatePrimarySpeciesConcentration, - dAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration, - dMobileAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration, - reactionRates, - dReactionRates_dlogPrimarySpeciesConcentration, - aggregateSpeciesRates, - dAggregateSpeciesRates_dlogPrimarySpeciesConcentration ); - - - for( int i = 0; i < numPrimarySpecies; ++i ) - { - r[i] = ( aggregatePrimarySpeciesConcentration[i] - aggregatePrimarySpeciesConcentration_n[i] ) - aggregateSpeciesRates[i] * dt; - for( int j = 0; j < numPrimarySpecies; ++j ) - { - J[i][j] = dAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration[i][j] - dAggregateSpeciesRates_dlogPrimarySpeciesConcentration[i][j] * dt; - } - } - }; + J[i][j] = dAggregatePrimarySpeciesConcentrations_dlogPrimarySpeciesConcentration[i][j] - dAggregateSpeciesRates_dlogPrimarySpeciesConcentration[i][j] * dt; + } + } + }; nonlinearSolvers::newtonRaphson< numPrimarySpecies >( logPrimarySpeciesConcentration, computeResidualAndJacobian ); diff --git a/src/uncrustify.cfg b/src/uncrustify.cfg index 6b3e1e0..0b0efda 100644 --- a/src/uncrustify.cfg +++ b/src/uncrustify.cfg @@ -1579,7 +1579,7 @@ pos_constr_colon = ignore # ignore/join/lead/lead_break/lead_fo # # Try to limit code width to N number of columns -code_width = 200 # unsigned number +code_width = 300 # unsigned number # Whether to fully split long 'for' statements at semi-colons. ls_for_split_full = false # false/true