// validate stencil-independent values
void
CheckOptions( const OptionParser& opts )
{
// check matrix dimensions - must be 2d, must be positive
std::vector<long long> arrayDims = opts.getOptionVecInt( "customSize" );
if( arrayDims.size() != 2 )
{
throw InvalidArgValue( "overall size must have two dimensions" );
}
if( (arrayDims[0] < 0) || (arrayDims[1] < 0) )
{
throw InvalidArgValue( "each size dimension must be positive" );
}
// validation error threshold must be positive
float valThreshold = opts.getOptionFloat( "val-threshold" );
if( valThreshold <= 0.0f )
{
throw InvalidArgValue( "validation threshold must be positive" );
}
// number of validation errors to print must be non-negative
int nErrsToPrint = opts.getOptionInt( "val-print-limit" );
if( nErrsToPrint < 0 )
{
throw InvalidArgValue( "number of validation errors to print must be non-negative" );
}
int nWarmupPasses = opts.getOptionInt( "warmupPasses" );
if( nWarmupPasses < 0 )
{
throw InvalidArgValue( "number of warmup passes must be non-negative" );
}
}
void
CommonMICStencilFactory<T>::CheckOptions( const OptionParser& opts ) const
{
// let base class check its options first
StencilFactory<T>::CheckOptions( opts );
// check our options
std::vector<long long> arrayDims = opts.getOptionVecInt( "customSize" );
assert( arrayDims.size() == 2 );
// If both of these are zero, we're using a non-custom size, skip this test
if (arrayDims[0] == 0 && arrayDims[0] == 0)
{
return;
}
size_t gRows = (size_t)arrayDims[0];
size_t gCols = (size_t)arrayDims[1];
size_t lRows = LROWS;
size_t lCols = LCOLS;
// verify that local dimensions evenly divide global dimensions
if( ((gRows % lRows) != 0) || (lRows > gRows) )
{
throw InvalidArgValue( "number of rows must be even multiple of lsize rows" );
}
if( ((gCols % lCols) != 0) || (lCols > gCols) )
{
throw InvalidArgValue( "number of columns must be even multiple of lsize columns" );
}
// TODO ensure local dims are smaller than CUDA implementation limits
}
void
MPICUDAStencilFactory<T>::CheckOptions( const OptionParser& opts ) const
{
// let base class check its options first
CommonCUDAStencilFactory<T>::CheckOptions( opts );
// check our options
std::vector<long long> shDims = opts.getOptionVecInt( "lsize" );
std::vector<long long> arrayDims = opts.getOptionVecInt( "customSize" );
if( arrayDims[0] == 0 )
{
// custom size was not specified - we are using a standard size
int sizeClass = opts.getOptionInt("size");
arrayDims = StencilFactory<T>::GetStandardProblemSize( sizeClass );
}
assert( shDims.size() == 2 );
assert( arrayDims.size() == 2 );
size_t gRows = (size_t)arrayDims[0];
size_t gCols = (size_t)arrayDims[1];
size_t lRows = shDims[0];
size_t lCols = shDims[1];
unsigned int haloWidth = (unsigned int)opts.getOptionInt( "iters-per-exchange" );
// verify that MPI halo width will result in a matrix being passed
// to the kernel that also has its global size as a multiple of
// the local work size
//
// Because the MPI halo width is arbitrary, and the kernel halo width
// is always 1, we have to ensure that:
// ((size + 2*halo) - 2) % lsize == 0
if( (((gRows + 2*haloWidth) - 2) % lRows) != 0 )
{
throw InvalidArgValue( "rows including halo must be even multiple of lsize (e.g., lsize rows evenly divides ((rows + 2*halo) - 2) )" );
}
if( (((gCols + 2*haloWidth) - 2) % lCols) != 0 )
{
throw InvalidArgValue( "columns including halo must be even multiple of lsize (e.g., lsize cols evenly divides ((cols + 2*halo) - 2) )" );
}
}
std::vector<long long>
StencilFactory<T>::GetStandardProblemSize( int sizeClass )
{
const int probSizes[4] = { 768, 1408, 2048, 4096 };
if (!(sizeClass >= 0 && sizeClass < 5))
{
throw InvalidArgValue( "Size class must be between 1-4" );
}
std::vector<long long> ret( 2, probSizes[sizeClass - 1] );
return ret;
}
void
StencilFactory<T>::CheckOptions( const OptionParser& options ) const
{
// number of iterations must be positive
unsigned int nIters = (unsigned int)options.getOptionInt( "num-iters" );
if( nIters == 0 )
{
throw InvalidArgValue( "number of iterations must be positive" );
}
// no restrictions on weight values, just that we have them
}
void
DoTest( const char* timerDesc, ResultDatabase& resultDB, OptionParser& opts )
{
StencilFactory<T>* stdStencilFactory = NULL;
Stencil<T>* stdStencil = NULL;
StencilFactory<T>* testStencilFactory = NULL;
Stencil<T>* testStencil = NULL;
try
{
#if defined(PARALLEL)
stdStencilFactory = new MPIHostStencilFactory<T>;
testStencilFactory = new MPICUDAStencilFactory<T>;
#else
stdStencilFactory = new HostStencilFactory<T>;
testStencilFactory = new CUDAStencilFactory<T>;
#endif // defined(PARALLEL)
assert( (stdStencilFactory != NULL) && (testStencilFactory != NULL) );
// do a sanity check on option values
CheckOptions( opts );
stdStencilFactory->CheckOptions( opts );
testStencilFactory->CheckOptions( opts );
// extract and validate options
std::vector<long long> arrayDims = opts.getOptionVecInt( "customSize" );
if( arrayDims.size() != 2 )
{
cerr << "Dim size: " << arrayDims.size() << "\n";
throw InvalidArgValue( "all overall dimensions must be positive" );
}
if (arrayDims[0] == 0) // User has not specified a custom size
{
int sizeClass = opts.getOptionInt("size");
arrayDims = StencilFactory<T>::GetStandardProblemSize( sizeClass );
}
long int seed = (long)opts.getOptionInt( "seed" );
bool beVerbose = opts.getOptionBool( "verbose" );
unsigned int nIters = (unsigned int)opts.getOptionInt( "num-iters" );
double valErrThreshold = (double)opts.getOptionFloat( "val-threshold" );
unsigned int nValErrsToPrint = (unsigned int)opts.getOptionInt( "val-print-limit" );
#if defined(PARALLEL)
unsigned int haloWidth = (unsigned int)opts.getOptionInt( "iters-per-exchange" );
#else
unsigned int haloWidth = 1;
#endif // defined(PARALLEL)
float haloVal = (float)opts.getOptionFloat( "haloVal" );
// build a description of this experiment
std::vector<long long> lDims = opts.getOptionVecInt( "lsize" );
assert( lDims.size() == 2 );
std::ostringstream experimentDescriptionStr;
experimentDescriptionStr
<< nIters << ':'
<< arrayDims[0] << 'x' << arrayDims[1] << ':'
<< lDims[0] << 'x' << lDims[1];
unsigned int nPasses = (unsigned int)opts.getOptionInt( "passes" );
unsigned int nWarmupPasses = (unsigned int)opts.getOptionInt( "warmupPasses" );
// compute the expected result on the host
// or read it from a pre-existing file
std::string matrixFilenameBase = (std::string)opts.getOptionString( "expMatrixFile" );
#if defined(PARALLEL)
int cwrank;
MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
if( cwrank == 0 )
{
#endif // defined(PARALLEL)
if( !matrixFilenameBase.empty() )
{
std::cout << "\nReading expected stencil operation result from file for later comparison with CUDA output\n"
<< std::endl;
}
else
{
std::cout << "\nPerforming stencil operation on host for later comparison with CUDA output\n"
<< "Depending on host capabilities, this may take a while."
<< std::endl;
}
#if defined(PARALLEL)
}
#endif // defined(PARALLEL)
Matrix2D<T> expected( arrayDims[0] + 2*haloWidth,
arrayDims[1] + 2*haloWidth );
Initialize<T> init( seed, haloWidth, haloVal );
bool haveExpectedData = false;
if( ! matrixFilenameBase.empty() )
{
bool readOK = ReadMatrixFromFile( expected, GetMatrixFileName<T>( matrixFilenameBase ) );
if( readOK )
{
if( (expected.GetNumRows() != arrayDims[0] + 2*haloWidth) ||
(expected.GetNumColumns() != arrayDims[1] + 2*haloWidth) )
{
std::cerr << "The matrix read from file \'"
<< GetMatrixFileName<T>( matrixFilenameBase )
<< "\' does not match the matrix size specified on the command line.\n";
expected.Reset( arrayDims[0] + 2*haloWidth, arrayDims[1] + 2*haloWidth );
}
else
{
haveExpectedData = true;
}
}
if( !haveExpectedData )
{
std::cout << "\nSince we could not read the expected matrix values,\nperforming stencil operation on host for later comparison with CUDA output.\n"
<< "Depending on host capabilities, this may take a while."
<< std::endl;
}
}
if( !haveExpectedData )
{
init( expected );
haveExpectedData = true;
if( beVerbose )
{
std::cout << "initial state:\n" << expected << std::endl;
}
stdStencil = stdStencilFactory->BuildStencil( opts );
(*stdStencil)( expected, nIters );
}
if( beVerbose )
{
std::cout << "expected result:\n" << expected << std::endl;
}
// determine whether we are to save the expected matrix values to a file
// to speed up future runs
matrixFilenameBase = (std::string)opts.getOptionString( "saveExpMatrixFile" );
if( !matrixFilenameBase.empty() )
{
SaveMatrixToFile( expected, GetMatrixFileName<T>( matrixFilenameBase ) );
}
assert( haveExpectedData );
// compute the result on the CUDA device
Matrix2D<T> data( arrayDims[0] + 2*haloWidth,
arrayDims[1] + 2*haloWidth );
Stencil<T>* testStencil = testStencilFactory->BuildStencil( opts );
// Compute the number of floating point operations we will perform.
//
// Note: in the truly-parallel case, we count flops for redundant
// work due to the need for a halo.
// But we do not add to the count for the local 1-wide halo since
// we aren't computing new values for those items.
unsigned long npts = (arrayDims[0] + 2*haloWidth - 2) *
(arrayDims[1] + 2*haloWidth - 2);
#if defined(PARALLEL)
MPICUDAStencil<T>* mpiTestStencil = static_cast<MPICUDAStencil<T>*>( testStencil );
assert( mpiTestStencil != NULL );
int participating = mpiTestStencil->ParticipatingInProgram() ? 1 : 0;
int numParticipating = 0;
MPI_Allreduce( &participating, // src
&numParticipating, // dest
1, // count
MPI_INT, // type
MPI_SUM, // op
MPI_COMM_WORLD ); // communicator
npts *= numParticipating;
#endif // defined(PARALLEL)
// In our 9-point stencil, there are 11 floating point operations
// per point (3 multiplies and 11 adds):
//
// newval = weight_center * centerval +
// weight_cardinal * (northval + southval + eastval + westval) +
// weight_diagnoal * (neval + nwval + seval + swval)
//
// we do this stencil operation 'nIters' times
unsigned long nflops = npts * 11 * nIters;
#if defined(PARALLEL)
if( cwrank == 0 )
{
#endif // defined(PARALLEL)
std::cout << "Performing " << nWarmupPasses << " warmup passes...";
#if defined(PARALLEL)
}
#endif // defined(PARALLEL)
for( unsigned int pass = 0; pass < nWarmupPasses; pass++ )
{
init(data);
(*testStencil)( data, nIters );
}
#if defined(PARALLEL)
if( cwrank == 0 )
{
#endif // defined(PARALLEL)
std::cout << "done." << std::endl;
#if defined(PARALLEL)
}
#endif // defined(PARALLEL)
#if defined(PARALLEL)
MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
if( cwrank == 0 )
{
#endif // defined(PARALLEL)
std::cout << "\nPerforming stencil operation on chosen device, "
<< nPasses << " passes.\n"
<< "Depending on chosen device, this may take a while."
<< std::endl;
#if defined(PARALLEL)
}
#endif // defined(PARALLEL)
#if !defined(PARALLEL)
std::cout << "At the end of each pass the number of validation\nerrors observed will be printed to the standard output."
<< std::endl;
#endif // !defined(PARALLEL)
for( unsigned int pass = 0; pass < nPasses; pass++ )
{
#if !defined(PARALLEL)
std::cout << "pass " << pass << ": ";
#endif // !defined(PARALLEL)
init( data );
int timerHandle = Timer::Start();
(*testStencil)( data, nIters );
double elapsedTime = Timer::Stop( timerHandle, "CUDA stencil" );
// find and report the computation rate
double gflops = (nflops / elapsedTime) / 1e9;
resultDB.AddResult( timerDesc,
experimentDescriptionStr.str(),
"GFLOPS",
gflops );
if( beVerbose )
{
std::cout << "observed result, pass " << pass << ":\n"
<< data
<< std::endl;
}
// validate the result
#if defined(PARALLEL)
StencilValidater<T>* validater = new MPIStencilValidater<T>;
#else
StencilValidater<T>* validater = new SerialStencilValidater<T>;
#endif // defined(PARALLEL)
validater->ValidateResult( expected,
data,
valErrThreshold,
nValErrsToPrint );
}
}
catch( ... )
{
// clean up - abnormal termination
// wish we didn't have to do this, but C++ exceptions do not
// support a try-catch-finally approach
delete stdStencil;
delete stdStencilFactory;
delete testStencil;
delete testStencilFactory;
throw;
}
// clean up - normal termination
delete stdStencil;
delete stdStencilFactory;
delete testStencil;
delete testStencilFactory;
}
