Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 18 additions & 12 deletions src/common/pmpl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down Expand Up @@ -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 );
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ void testMoMasAllEquilibriumHelper()
targetAggregatePrimarySpeciesConcentration,
logInitialPrimarySpeciesConcentration,
logPrimarySpeciesConcentrationCopy );
});
} );

double const expectedPrimarySpeciesConcentrations[numPrimarySpecies] =
{
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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] =
{
Expand Down
11 changes: 4 additions & 7 deletions src/reactions/massActions/MassActions.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 ){} );
}


Expand Down Expand Up @@ -154,8 +154,7 @@ void calculateAggregatePrimaryConcentrationsWrtLogC( PARAMS_DATA const & params,
for( int k=0; k<numPrimarySpecies; ++k )
{
REAL_TYPE const dSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentration = params.stoichiometricMatrix( j, k+numSecondarySpecies ) * secondarySpeciesConcentrations_j;
dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j,
i+numSecondarySpecies ) *
dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j, i+numSecondarySpecies ) *
dSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentration;
}
}
Expand Down Expand Up @@ -248,12 +247,10 @@ void calculateTotalAndMobileAggregatePrimaryConcentrationsWrtLogC( PARAMS_DATA c
for( int k=0; k<numPrimarySpecies; ++k )
{
REAL_TYPE const dSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentration = params.stoichiometricMatrix( j, k+numSecondarySpecies ) * secondarySpeciesConcentrations_j;
dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j,
i+numSecondarySpecies ) *
dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j, i+numSecondarySpecies ) *
dSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentration;

dMobileAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j,
i+numSecondarySpecies ) *
dMobileAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations( i, k ) += params.stoichiometricMatrix( j, i+numSecondarySpecies ) *
dSecondarySpeciesConcentrations_dLogPrimarySpeciesConcentration *
params.mobileSecondarySpeciesFlag( j );
}
Expand Down
33 changes: 16 additions & 17 deletions src/reactions/massActions/unitTests/testMassActions.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -45,13 +45,13 @@ void test_calculateLogSecondarySpeciesConcentration_helper()
static constexpr int numPrimarySpecies = carbonateSystemAllEquilibrium.numPrimarySpecies();
static constexpr int numSecondarySpecies = carbonateSystemAllEquilibrium.numSecondarySpecies();

CalculateLogSecondarySpeciesConcentrationData<numPrimarySpecies, numSecondarySpecies> 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 );

Expand All @@ -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 );
} );



Expand Down Expand Up @@ -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};
Expand All @@ -144,15 +143,15 @@ void testcalculateAggregatePrimaryConcentrationsWrtLogCHelper()
{
static constexpr int numPrimarySpecies = carbonateSystemAllEquilibrium.numPrimarySpecies();

CalculateAggregatePrimaryConcentrationsWrtLogCHelperData<numPrimarySpecies> data;
CalculateAggregatePrimaryConcentrationsWrtLogCHelperData< numPrimarySpecies > data;

pmpl::genericKernelWrapper( 1, &data, [] HPCREACT_DEVICE ( auto * const dataCopy )
{
calculateAggregatePrimaryConcentrationsWrtLogC< double, int, int >( carbonateSystemAllEquilibrium.equilibriumReactionsParameters(),
dataCopy->primarySpeciesSolution,
dataCopy->aggregatePrimarySpeciesConcentration,
dataCopy->dAggregatePrimarySpeciesConcentrationsDerivatives_dLogPrimarySpeciesConcentrations );
});
} );

double const expectedAggregatePrimarySpeciesConcentration[numPrimarySpecies] =
{
Expand Down
14 changes: 7 additions & 7 deletions src/reactions/reactionsSystems/Parameters.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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;
Expand Down Expand Up @@ -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;
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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 };
Expand Down Expand Up @@ -65,23 +65,23 @@ void computeResidualAndJacobianTest( PARAMS_DATA const & params,

double const temperature = 298.15;

ComputeResidualAndJacobianTestData<numReactions, numSpecies> 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<numReactions; ++r )
Expand Down Expand Up @@ -114,7 +114,7 @@ struct TestEnforceEquilibriumData
{
/// The initial species concentrations
double speciesConcentration0[numSpecies];

/// The final species concentrations
double speciesConcentration[numSpecies];
};
Expand All @@ -134,19 +134,19 @@ void testEnforceEquilibrium( PARAMS_DATA const & params,

double const temperature = 298.15;

TestEnforceEquilibriumData<numSpecies> 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; r<numSpecies; ++r )
{
Expand Down
Loading