예제 #1
0
// 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" );
    }
}
예제 #2
0
// ****************************************************************************
// Function: RunBenchmark
//
// Purpose:
//   Measures the floating point capability of the device for a variety of 
//   combinations of arithmetic operations.
//
// Arguments:
//   op: the options parser / parameter database
//
// Returns:  nothing
//
// Programmer: Zhi Ying([email protected])
//             Jun Jin([email protected])
//
// Creation: May 23, 2011
//
// Modifications:
// 12/12/12 - Kyle Spafford - Code style and minor integration updates
//
// ****************************************************************************
void RunBenchmark(OptionParser &op, ResultDatabase &resultDB)
{
    const bool verbose = op.getOptionBool("verbose");
    // Quiet == no progress bar.
    const bool quiet   = op.getOptionBool("quiet");
    const unsigned int passes = op.getOptionInt("passes");
    const int micdev = op.getOptionInt("target");

    double repeatF = 3;
    cout << "Adjust repeat factor = " << repeatF << "\n";

    // Initialize progress bar
    int totalRuns = 16*passes*2;
    ProgressBar pb(totalRuns);
    if (!verbose && !quiet) 
    {
        pb.Show(stdout);
    }

    RunTest<float>(resultDB, passes, verbose, quiet,
                   repeatF, pb, "-SP", micdev);
    RunTest<double>(resultDB, passes, verbose, quiet,
                    repeatF, pb, "-DP", micdev);

    if (!verbose) cout << endl;
}
예제 #3
0
// ****************************************************************************
// Function: GPUSetup
//
// Purpose:
//  do the necessary OpenCL setup for GPU part of the test
//
// Arguments:
//   op: the options parser / parameter database
//   mympirank: for printing errors in case of failure
//   mynoderank: this is typically the device ID (the mapping done in main)
//
// Returns: success/failure
//
// Creation: 2009
//
// Modifications:
//
// ****************************************************************************
//
int GPUSetup(OptionParser &op, int mympirank, int mynoderank)
{
    addBenchmarkSpecOptions(op);
    
    if (op.getOptionBool("infoDevices"))
    {
        OpenCLNodePlatformContainer ndc1;
        ndc1.Print (cout);
        return (0);
    }
    
    // The device option supports specifying more than one device
    int platform = op.getOptionInt("platform");
    int deviceIdx = mynoderank;
    if( deviceIdx >= op.getOptionVecInt( "device" ).size() )
    {
        std::ostringstream estr;
        estr << "Warning: not enough devices specified with --device flag for task "
            << mympirank
            << " ( node rank " << mynoderank
            << ") to claim its own device; forcing to use first device ";
        std::cerr << estr.str() << std::endl;
        deviceIdx = 0;
    }
    int device = op.getOptionVecInt("device")[deviceIdx];

    // Initialization
    _mpicontention_ocldev = new cl::Device( ListDevicesAndGetDevice(platform, device) );
    std::vector<cl::Device> ctxDevices;
    ctxDevices.push_back( *_mpicontention_ocldev );
    _mpicontention_ocldriver_ctx   = new cl::Context( ctxDevices );
    _mpicontention_ocldriver_queue = new cl::CommandQueue( *_mpicontention_ocldriver_ctx, *_mpicontention_ocldev, CL_QUEUE_PROFILING_ENABLE );
    _mpicontention_gpuop = op;
    return 0;
}
예제 #4
0
파일: FFT.cpp 프로젝트: vetter/shoc
void
RunBenchmark(ResultDatabase &resultDB, OptionParser &op)
{
    // Test to see if this device supports double precision
    cudaGetDevice(&fftDevice);
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, fftDevice);
    bool has_dp = (deviceProp.major == 1 && deviceProp.minor >= 3) ||
        (deviceProp.major >= 2);

    cout << "Running single precision test" << endl;
    runTest<float2>("SP-FFT", resultDB, op);
    if (has_dp) {
        cout << "Running double precision test" << endl;
        runTest<double2>("DP-FFT", resultDB, op);
    }
    else
    {
        cout << "Skipping double precision test" << endl;
        char atts[32] = "DP_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (int k=0; k<passes; k++)
        {
            resultDB.AddResult("DP-FFT" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("DP-FFT_PCIe" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("DP-FFT_Parity" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("DP-FFT-INV" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("DP-FFT-INV_PCIe" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("DP-FFT-INV_Parity" , atts, "GB/s", FLT_MAX);
        }
    }
}
예제 #5
0
파일: fft.cpp 프로젝트: Sable/Ostrich
void dump2D(OptionParser& op)
{
	int i;
	void* work, *temp;
	T2* source, * result;
	unsigned long bytes = 0;

	int probSizes[7] = { 128, 256, 512, 1024, 2048, 4096, 8192};
	int sizeIndex = op.getOptionInt("pts1")-1;
	int sizeIndey = op.getOptionInt("pts2")-1;
	if (sizeIndex < 0 || sizeIndex >= 7) {
		cerr << "Invalid size index specified\n";
		exit(-1);
	}
	if (sizeIndey < 0 || sizeIndey >= 7) {
		cerr << "Invalid size index specified\n";
		exit(-1);
	}

	int FFTN1=probSizes[sizeIndex],FFTN2=probSizes[sizeIndey];
	//int FFTN1=8192,FFTN2=512;
	unsigned long used_bytes = FFTN1*FFTN2*sizeof(T2);

	bool do_dp = dp<T2>();
	init2(op, do_dp, FFTN1, FFTN2);

	int n_ffts = 1;
	double N = FFTN1*FFTN2;


	// allocate host and device memory
	allocHostBuffer((void**)&source, used_bytes);
	allocHostBuffer((void**)&result, used_bytes);

	// init host memory...
	for (i = 0; i < N; i++) {
		source[i].x = (rand()/(float)RAND_MAX)*2-1;
		source[i].y = (rand()/(float)RAND_MAX)*2-1;
	}

	// alloc device memory
	allocDeviceBuffer(&work, used_bytes);
	allocDeviceBuffer(&temp, used_bytes);

	copyToDevice(work, source, used_bytes);

	forward2(work, temp, n_ffts, FFTN1, FFTN2);
	copyFromDevice(result, work, used_bytes);

#ifdef PRINT_RESULT
	for (i = 0; i < N; i++) {
		fprintf(stdout, "data[%d] (%g, %g) \n",i, result[i].x, result[i].y);
	}
#endif
	freeDeviceBuffer(work);
	freeDeviceBuffer(temp);
	freeHostBuffer(source);
	freeHostBuffer(result);
}
예제 #6
0
파일: tpScan.cpp 프로젝트: dylanzika/shoc
void
RunBenchmark(cl::Device& devcpp,
             cl::Context& ctxcpp,
             cl::CommandQueue& queuecpp,
             ResultDatabase &resultDB,
             OptionParser &op)
{
    // Convert from C++ bindings to C bindings
    // TODO propagate use of C++ bindings
    cl_device_id dev = devcpp();
    cl_context ctx = ctxcpp();
    cl_command_queue queue = queuecpp();

    // Collect basic MPI information
    int size, rank;
    MPI_Comm_size(MPI_COMM_WORLD, &size);
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);

    // Always run single precision test
    // OpenCL doesn't support templated kernels, so we have to use macros
    string spMacros = "-DSINGLE_PRECISION";
    runTest<float>
        ("TPScan-SP", dev, ctx, queue, resultDB, op, spMacros);

    // If double precision is supported, run the DP test
    if (checkExtension(dev, "cl_khr_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DK_DOUBLE_PRECISION ";
        runTest<double>
        ("TPScan-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else if (checkExtension(dev, "cl_amd_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DAMD_DOUBLE_PRECISION ";
        runTest<double>
        ("TPScan-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else
    {
        char atts[1024] = "DP_Not_Supported";
        cout << "Warning, rank " << rank << "'s device does not support DP\n";
        // ResultDB requires every rank to report something. If this rank
        // doesn't support DP, submit FLT_MAX (this is handled as no result by
        // ResultDB.
        int passes = op.getOptionInt("passes");
        for (int k = 0; k < passes; k++)
        {
            resultDB.AddResult("TPScan-DP-Kernel" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("TPScan-DP-Kernel+PCIe" , atts, "GB/s",
                FLT_MAX);
            resultDB.AddResult("TPScan-DP-MPI_ExScan" , atts, "GB/s",
                FLT_MAX);
            resultDB.AddResult("TPScan-DP-Overall" , atts, "GB/s", FLT_MAX);
        }
    }
}
예제 #7
0
파일: fft.cpp 프로젝트: Sable/Ostrich
void dump1D(OptionParser& op)
{
	int i;
	int fftn;
	void* work, *temp;
	T2* source, * result;
	unsigned long bytes = 0;

	int probSizes[7] = { 128, 256, 512, 1024, 2048, 4096, 8192 };
	int sizeIndex = op.getOptionInt("pts")-1;
	if (sizeIndex < 0 || sizeIndex >= 7) {
		cerr << "Invalid size index specified\n";
		exit(-1);
	}
	fftn = probSizes[sizeIndex];

	// Convert to MB
	unsigned long used_bytes = fftn * sizeof(T2);

	bool do_dp = dp<T2>();
	init(op, do_dp, fftn);

	// now determine how much available memory will be used
	//int half_n_ffts = bytes / (fftn*sizeof(T2)*2);
	int n_ffts = 1;
	double N = fftn;

	fprintf(stdout, "used_bytes=%lu, N=%g\n", used_bytes, N);

	// allocate host and device memory
	allocHostBuffer((void**)&source, used_bytes);
	allocHostBuffer((void**)&result, used_bytes);

	// init host memory...
	for (i = 0; i < N; i++) {
		source[i].x = (rand()/(float)RAND_MAX)*2-1;
		source[i].y = (rand()/(float)RAND_MAX)*2-1;
	}

	// alloc device memory
	allocDeviceBuffer(&work, used_bytes);
	allocDeviceBuffer(&temp, used_bytes);
	copyToDevice(work, source, used_bytes);


	forward(work, temp, n_ffts, fftn);

	copyFromDevice(result, work, used_bytes);
#ifdef PRINT_RESULT
	for (i = 0; i < N; i++) {
		fprintf(stdout, "data[%d] (%g, %g)\n", i, result[i].x, result[i].y);
	}
#endif
	freeDeviceBuffer(work);
	freeDeviceBuffer(temp);
	freeHostBuffer(source);
	freeHostBuffer(result);
}
예제 #8
0
void
RunBenchmark(OptionParser& opts, ResultDatabase& resultDB )
{
    int device;

#if defined(PARALLEL)
    int cwrank;
    MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
#endif // defined(PARALLEL)
#if defined(PARALLEL)
    if( cwrank == 0 )
    {
#endif // defined(PARALLEL)
        std::cout << "Running single precision test" << std::endl;
#if defined(PARALLEL)
    }
#endif // defined(PARALLEL)
    //omp_set_num_threads(124);
    DoTest<float>( "SP_Sten2D", resultDB, opts );

    // check if we can run double precision tests
    if( //deviceProps.major == 1) && (deviceProps.minor >= 3)) ||
        //eviceProps.major >= 2))
        1)
    {
#if defined(PARALLEL)
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "DP supported\n" << std::endl;
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
	//omp_set_num_threads(93);
        DoTest<double>( "DP_Sten2D", resultDB, opts );
    }
    else
    {
#if defined(PARALLEL)
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "Double precision not supported - skipping" << std::endl;
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
        // resultDB requires neg entry for every possible result
        int nPasses = (int)opts.getOptionInt( "passes" );
        for( int p = 0; p < nPasses; p++ )
        {
            resultDB.AddResult( (const char*)"DP_Sten2D", "N/A", "s", FLT_MAX);
        }
    }
	
}
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) )" );
    }
}
예제 #10
0
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
}
예제 #11
0
파일: FFT.cpp 프로젝트: Poojachitral/shoc
static void
fillResultDB(const string& name, const string& reason, OptionParser &op, 
             ResultDatabase& resultDB)
{
    // resultDB requires neg entry for every possible result
    int passes = op.getOptionInt("passes");
    for (int k=0; k<passes; k++) {
        resultDB.AddResult(name , reason, "GB/s", FLT_MAX);
        resultDB.AddResult(name+"_PCIe" , reason, "GB/s", FLT_MAX);
        resultDB.AddResult(name+"_Parity" , reason, "GB/s", FLT_MAX);
        resultDB.AddResult(name+"-INV" , reason, "GB/s", FLT_MAX);
        resultDB.AddResult(name+"-INV_PCIe" , reason, "GB/s", FLT_MAX);
        resultDB.AddResult(name+"-INV_Parity" , reason, "GB/s", FLT_MAX);
    }
}
예제 #12
0
void
RunBenchmark(cl::Device& devcpp,
                  cl::Context& ctxcpp,
                  cl::CommandQueue& queuecpp,
                  ResultDatabase &resultDB,
                  OptionParser &op)
{
    // convert from C++ bindings to C bindings
    // TODO propagate use of C++ bindings
    cl_device_id dev = devcpp();
    cl_context ctx = ctxcpp();
    cl_command_queue queue = queuecpp();

    // Always run single precision test
    // OpenCL doesn't support templated kernels, so we have to use macros
    runTest<float>("SGEMM", dev, ctx, queue, resultDB, op,
            "-DSINGLE_PRECISION");

    // If double precision is supported, run the DP test
    if (checkExtension(dev, "cl_khr_fp64"))
    {
        cout << "DP Supported\n";
        runTest<double>("DGEMM", dev, ctx, queue, resultDB, op,
                "-DK_DOUBLE_PRECISION ");
    }
    else if (checkExtension(dev, "cl_amd_fp64"))
    {
        cout << "DP Supported\n";
        runTest<double>("DGEMM", dev, ctx, queue, resultDB, op,
                "-DAMD_DOUBLE_PRECISION ");
    }
    else
    {
        cout << "DP Not Supported\n";
        char atts[1024] = "DP_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (; passes > 0; --passes) {
            for (int i = 0; i < 2; i++) {
                const char transb = i ? 'T' : 'N';
                string testName="DGEMM";
                resultDB.AddResult(testName+"-"+transb, atts, "GFlops", FLT_MAX);
                resultDB.AddResult(testName+"-"+transb+"_PCIe", atts, "GFlops", FLT_MAX);
                resultDB.AddResult(testName+"-"+transb+"_Parity", atts, "N", FLT_MAX);
            }
        }
    }
}
예제 #13
0
파일: S3D.cpp 프로젝트: dylanzika/shoc
void
RunBenchmark(cl::Device& devcpp, cl::Context& ctxcpp,
             cl::CommandQueue& queuecpp,
             ResultDatabase &resultDB, OptionParser &op)
{
    // convert from C++ bindings to C bindings
    // TODO propagate use of C++ bindings
    cl_device_id dev = devcpp();
    cl_context ctx = ctxcpp();
    cl_command_queue queue = queuecpp();

    // Always run single precision test
    // OpenCL doesn't support templated kernels, so we have to use macros
    string spMacros = "-DSINGLE_PRECISION ";
    RunTest<float>("S3D-SP", dev, ctx, queue, resultDB, op, spMacros);

    // If double precision is supported, run the DP test
    if (checkExtension(dev, "cl_khr_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DK_DOUBLE_PRECISION ";
        RunTest<double>
        ("S3D-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else if (checkExtension(dev, "cl_amd_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DAMD_DOUBLE_PRECISION ";
        RunTest<double>
        ("S3D-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else
    {
        cout << "DP Not Supported\n";
        char atts[1024] = "DP_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (int k = 0; k < passes; k++) {
            resultDB.AddResult("S3D-DP" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("S3D-DP_PCIe" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("S3D-DP_Parity" , atts, "GB/s", FLT_MAX);
        }
    }
}
예제 #14
0
void
RunBenchmark(cl_device_id dev,
                  cl_context ctx,
                  cl_command_queue queue,
                  ResultDatabase &resultDB,
                  OptionParser &op)
{
    // Always run single precision test
    // OpenCL doesn't support templated kernels, so we have to use macros
    string spMacros = "-DSINGLE_PRECISION";
    runTest<float, float4, float4>
        ("MD-LJ", dev, ctx, queue, resultDB, op, spMacros);

    // If double precision is supported, run the DP test
    if (checkExtension(dev, "cl_khr_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DK_DOUBLE_PRECISION ";
        runTest<double, double4, double4>
                ("MD-LJ-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else if (checkExtension(dev, "cl_amd_fp64"))
    {
        cout << "DP Supported\n";
        string dpMacros = "-DAMD_DOUBLE_PRECISION ";
        runTest<double, double4, double4>
        ("MD-LJ-DP", dev, ctx, queue, resultDB, op, dpMacros);
    }
    else
    {
        cout << "DP Not Supported\n";
        char atts[32] = "DP_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (int i = 0; i < passes; i++) {
            resultDB.AddResult("MD-LJ-DP" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("MD-LJ-DP_PCIe" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult("MD-LJ-DP-Bandwidth", atts, "GB/s", FLT_MAX);
            resultDB.AddResult("MD-LJ-DP-Bandwidth_PCIe", atts, "GB/s", FLT_MAX);
            resultDB.AddResult("MD-LJ-DP_Parity" , atts, "GB/s", FLT_MAX);
        }
    }
}
예제 #15
0
// ****************************************************************************
// Function: RunBenchmark
//
// Purpose: 
//   Runs the stablity test. The algorithm for the parallel
//   version of the test, which enables testing of an entire GPU
//   cluster at the same time, is as follows. Each participating node
//   first allocates its data, while node zero additionally determines
//   start and finish times based on a user input parameter. All nodes
//   then enter the outermost loop, copying fresh data from the CPU
//   before entering the core of the test. In the core, each node
//   performs a loop consisting of the forward kernel, a potential
//   check, and then the inverse kernel. After performing a configurable
//   number of forward/inverse iterations, along with a configurable
//   number of checks, each node sends the number of failures it
//   encountered to node zero. Node zero collects and reports the error
//   counts, determines whether the test has run its course, and
//   broadcasts the decision. If the decision is to proceed, each node
//   begins the next iteration of the outer loop, copying fresh data and
//   then performing the kernels and checks of the core loop.
//
// Arguments:
//   resultDB: the benchmark stores its results in this ResultDatabase
//   op: the options parser / parameter database
//
// Returns:  nothing
//
// Programmer: Collin McCurdy
// Creation: September 08, 2009
//
// Modifications:
//
// ****************************************************************************
void RunBenchmark(ResultDatabase &resultDB, OptionParser& op)
{	
    int mpi_rank, mpi_size, node_rank;
    int i, j;
    float2* source, * result;
    void* work, * chk;

#ifdef PARALLEL
    MPI_Comm_size(MPI_COMM_WORLD, &mpi_size);
    MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);

    NodeInfo NI;
    node_rank = NI.nodeRank();

    cout << "MPI Task " << mpi_rank << " of " << mpi_size 
         << " (noderank=" << node_rank << ") starting....\n";
#else
    mpi_rank = 0;
    mpi_size = 1;
    node_rank = 0;
#endif
   
    // ensure chk buffer alloc succeeds before grabbing the
    // rest of available memory.
    allocDeviceBuffer(&chk, 1);
    unsigned long avail_bytes = findAvailBytes();
    // unsigned long avail_bytes = 1024*1024*1024-1;
    
    // now determine how much available memory will be used (subject
    // to CUDA's constraint on the maximum block dimension size)
    int blocks = avail_bytes / (512*sizeof(float2));
    int slices = 1;
    while (blocks/slices > 65535) {
        slices *= 2;
    }
    int half_n_ffts = ((blocks/slices)*slices)/2;
    int n_ffts = half_n_ffts * 2;
    fprintf(stderr, "avail_bytes=%ld, blocks=%d, n_ffts=%d\n", 
            avail_bytes, blocks, n_ffts);

    int half_n_cmplx = half_n_ffts * 512;
    unsigned long used_bytes = half_n_cmplx * 2 * sizeof(float2);

    cout << mpi_rank << ": testing " 
         << used_bytes/((double)1024*1024) << " MBs\n";

    // allocate host memory
    source = (float2*)malloc(used_bytes);
    result = (float2*)malloc(used_bytes);

    // alloc device memory
    allocDeviceBuffer(&work, used_bytes);

    // alloc gather buffer
    int* recvbuf = (int*)malloc(mpi_size*sizeof(int));
    
    // compute start and finish times
    time_t start = time(NULL);
    time_t finish = start + (time_t)(op.getOptionInt("time")*60);
    struct tm start_tm, finish_tm;
    localtime_r(&start, &start_tm);
    localtime_r(&finish, &finish_tm);
    if (mpi_rank == 0) {
        printf("start = %s", asctime(&start_tm));
        printf("finish = %s", asctime(&finish_tm));
    }
    
    for (int iter = 0; ; iter++) {
        bool failed = false;
        int errorCount = 0, stop = 0;

        // (re-)init host memory...
        for (i = 0; i < half_n_cmplx; i++) {
            source[i].x = (rand()/(float)RAND_MAX)*2-1;
            source[i].y = (rand()/(float)RAND_MAX)*2-1;
            source[i+half_n_cmplx].x = source[i].x;
            source[i+half_n_cmplx].y = source[i].y;
        }

        // copy to device
        copyToDevice(work, source, used_bytes);
        copyToDevice(chk, &errorCount, 1);

        forward(work, n_ffts);
        if (check(work, chk, half_n_ffts, half_n_cmplx)) {
            fprintf(stderr, "First check failed...");
            failed = true;
        }

        if (!failed) {
            for (i = 1; i <= CHECKS; i++) {
                for (j = 1; j <= ITERS_PER_CHECK; j++) {
                    inverse(work, n_ffts);
                    forward(work, n_ffts);
                }

                if (check(work, chk, half_n_ffts, half_n_cmplx)) {
                    failed = true;
                    break;
                }
            }
        }

        // failing node is responsible for verifying failure, counting
        // errors and reporting count to node 0.
        if (failed) {
            fprintf(stderr, "Failure on node %d, iter %d:", mpi_rank, iter);
            
            // repeat check on CPU
            copyFromDevice(result, work, used_bytes);
            float2* result2 = result + half_n_cmplx;
            for (j = 0; j < half_n_cmplx; j++) {
                if (result[j].x != result2[j].x || 
                    result[j].y != result2[j].y) 
                {
                    errorCount++;
                }
            }
            if (!errorCount) {
                fprintf(stderr, "verification failed!\n");
            }
            else {
                fprintf(stderr, "%d errors\n", errorCount);
            }
        }

#ifdef PARALLEL
        MPI_Gather(&errorCount, 1, MPI_INT, recvbuf, 1, MPI_INT, 
                   0, MPI_COMM_WORLD);
#else
        recvbuf[0] = errorCount;
#endif

        // node 0 collects and reports error counts, determines
        // whether test has run its course, and broadcasts decision
        if (mpi_rank == 0) {
            time_t curtime = time(NULL);
            struct tm curtm;
            localtime_r(&curtime, &curtm);
            fprintf(stderr, "iter=%d: %s", iter, asctime(&curtm));
            
            for (i = 0; i < mpi_size; i++) {
                if (recvbuf[i]) {
                    fprintf(stderr, "--> %d failures on node %d\n", recvbuf[i], i);
                }
            }

            if (curtime > finish) {
                stop = 1;
            }
        }
        
#ifdef PARALLEL
        MPI_Bcast(&stop, 1, MPI_INT, 0, MPI_COMM_WORLD);
#endif

        resultDB.AddResult("Check", "", "Failures", errorCount);
        if (stop) break;
    }
    
    freeDeviceBuffer(work);
    freeDeviceBuffer(chk);

    free(source);
    free(result);
    free(recvbuf);
}
예제 #16
0
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;
}
예제 #17
0
void
RunBenchmark( ResultDatabase& resultDB, OptionParser& opts )
{
    int device;

#if defined(PARALLEL)
    int cwrank;
    MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
#endif // defined(PARALLEL)

    cudaGetDevice( &device );
    cudaDeviceProp deviceProps;
    cudaGetDeviceProperties( &deviceProps, device );

    // Configure to allocate performance-critical memory in
    // a programming model-specific way.
    Matrix2D<float>::SetAllocator( new CUDAPMSMemMgr<float> );

#if defined(PARALLEL)
    if( cwrank == 0 )
    {
#endif // defined(PARALLEL)
        std::cout << "Running single precision test" << std::endl;
#if defined(PARALLEL)
    }
#endif // defined(PARALLEL)
    DoTest<float>( "SP_Sten2D", resultDB, opts );

    // check if we can run double precision tests
    if( ((deviceProps.major == 1) && (deviceProps.minor >= 3)) ||
        (deviceProps.major >= 2) )
    {
        // Configure to allocate performance-critical memory in
        // a programming model-specific way.
        Matrix2D<double>::SetAllocator( new CUDAPMSMemMgr<double> );

#if defined(PARALLEL)
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "\n\nDP supported" << std::endl;
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
        DoTest<double>( "DP_Sten2D", resultDB, opts );
    }
    else
    {
#if defined(PARALLEL)
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "Double precision not supported - skipping" << std::endl;
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
        // resultDB requires neg entry for every possible result
        int nPasses = (int)opts.getOptionInt( "passes" );
        for( int p = 0; p < nPasses; p++ )
        {
            resultDB.AddResult( (const char*)"DP_Sten2D", "N/A", "GFLOPS", FLT_MAX );
        }
    }

    std::cout << "\n" << std::endl;
}
void
RunBenchmark( cl::Device& dev,
                cl::Context& ctx,
                cl::CommandQueue& queue,
                ResultDatabase& resultDB,
                OptionParser& op )
{
#if defined(PARALLEL)
    int cwrank;
#endif // defined(PARALLEL)

    // single precision
    DoTest<float>( "SP_Sten2D", dev, ctx, queue, resultDB, op, "-DSINGLE_PRECISION" );

    // double precision - might not be supported
    if( checkExtension( dev, "cl_khr_fp64" ))
    {
#if defined(PARALLEL)
        MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "\nDP supported\n";
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
        DoTest<double>( "DP_Sten2D", dev, ctx, queue, resultDB, op, "-DK_DOUBLE_PRECISION" );
    }
    else if( checkExtension( dev, "cl_amd_fp64" ))
    {
#if defined(PARALLEL)
        MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "\nDP supported\n";
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)
        DoTest<double>( "DP_Sten2D", dev, ctx, queue, resultDB, op, "-DAMD_DOUBLE_PRECISION" );
    }
    else
    {
#if defined(PARALLEL)
        MPI_Comm_rank( MPI_COMM_WORLD, &cwrank );
        if( cwrank == 0 )
        {
#endif // defined(PARALLEL)
            std::cout << "\nDP not supported\n";
#if defined(PARALLEL)
        }
#endif // defined(PARALLEL)

        // resultDB requires an entry for every possible result
        int nPasses = (int)op.getOptionInt( "passes" );
        for( unsigned int p = 0; p < nPasses; p++ )
        {
            resultDB.AddResult( (const char*)"DP_Sten2D",
                                "N/A",
                                "GFLOPS",
                                FLT_MAX );
        }
    }

    std::cout << '\n' << std::endl;
}
예제 #19
0
void runTest(const string& testName, cl_device_id dev, cl_context ctx,
        cl_command_queue queue, ResultDatabase& resultDB, OptionParser& op,
        const string& compileFlags)
{
    // Collect basic MPI information
    int mpi_size, mpi_rank;
    MPI_Comm_size(MPI_COMM_WORLD, &mpi_size);
    MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);

    int err;
    int waitForEvents = 1;

    // Program Setup
    cl_program prog = clCreateProgramWithSource(ctx, 1,
                            &cl_source_reduction, NULL, &err);
    CL_CHECK_ERROR(err);
    if (mpi_rank == 0)
    {
        cout << "Compiling reduction kernel." << endl;
    }

    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    CL_CHECK_ERROR(err);

    if (err != 0)
    {
        char log[5000];
        size_t retsize = 0;
        err =  clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG,
                5000*sizeof(char), log, &retsize);

        CL_CHECK_ERROR(err);
        cout << "Build error." << endl;
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

    // Extract out the kernels
    cl_kernel reduce = clCreateKernel(prog, "reduce", &err);
    CL_CHECK_ERROR(err);

    cl_kernel cpureduce = clCreateKernel(prog, "reduceNoLocal", &err);
    CL_CHECK_ERROR(err);

    size_t localWorkSize = 256;
    bool nolocal = false;
    if (getMaxWorkGroupSize(ctx, reduce) == 1) 
    {
        nolocal = true;
        localWorkSize = 1;
    }

    int probSizes[4] = { 1, 8, 64, 128 };

    int size = probSizes[op.getOptionInt("size")-1];
    size = (size * 1024 * 1024) / sizeof(T);

    unsigned int bytes = size * sizeof(T);

    // Allocate pinned host memory for input data
    cl_mem h_i = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_idata = (T*)clEnqueueMapBuffer(queue, h_i, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);

    // Initialize host memory
    if (mpi_rank == 0)
    {
        cout << "Initializing host memory." << endl;
    }

    for(int i=0; i<size; i++)
    {
        h_idata[i] = i % 2; //Fill with some pattern
    }

    // Allocate device memory for input data
    cl_mem d_idata = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes,
            NULL, &err);
    CL_CHECK_ERROR(err);

    int num_blocks;
    if (!nolocal)
    {
        num_blocks = 64;
    }
    else
    {
        num_blocks = 1; // NB: This should only be the case on Apple's CPU
                       // implementation, which is quite restrictive on
                       // work group sizes.
    }

    // Allocate host memory for output
    cl_mem h_o = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            sizeof(T)*num_blocks, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_odata = (T*)clEnqueueMapBuffer(queue, h_o, true,
            CL_MAP_READ | CL_MAP_WRITE, 0, sizeof(T) * num_blocks , 0, NULL, NULL,
            &err);
    CL_CHECK_ERROR(err);

    // Allocate device memory for output
    cl_mem d_odata = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
            num_blocks * sizeof(T), NULL, &err);
    CL_CHECK_ERROR(err);

    // Copy data to GPU
    Event evTransfer("PCIe Transfer");
    err = clEnqueueWriteBuffer(queue, d_idata, true, 0, bytes, h_idata,
            0, NULL, &evTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    evTransfer.FillTimingInfo();

    double inputTransfer = evTransfer.StartEndRuntime();

    err = clSetKernelArg(reduce, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 1, sizeof(cl_mem), (void*)&d_odata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 2,
            localWorkSize * sizeof(T), NULL);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 3, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);

    err = clSetKernelArg(cpureduce, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(cpureduce, 1, sizeof(cl_mem), (void*)&d_odata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(cpureduce, 2, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);

    size_t globalWorkSize;
    if (!nolocal)
    {
        globalWorkSize = localWorkSize * 64; // Use 64 work groups
    }
    else
    {
        globalWorkSize = 1;
    }

    int passes = op.getOptionInt("passes");
    int iters  = op.getOptionInt("iterations");

    if (mpi_rank == 0)
    {
        cout << "Running benchmark." << endl;
    }
    for (int k = 0; k < passes; k++)
    {
        // Synch processes at the start of each test.
        MPI_Barrier(MPI_COMM_WORLD); 
        
        double totalReduceTime = 0.0;
        Event evKernel("reduce kernel");
        for (int m = 0; m < iters; m++)
        {
            if (nolocal) 
            {
                err = clEnqueueNDRangeKernel(queue, cpureduce, 1, NULL,
                        &globalWorkSize, &localWorkSize, 0,
                        NULL, &evKernel.CLEvent());
            }
            else 
            {
                err = clEnqueueNDRangeKernel(queue, reduce, 1, NULL,
                        &globalWorkSize, &localWorkSize, 0,
                        NULL, &evKernel.CLEvent());
            }
            CL_CHECK_ERROR(err);
            err = clFinish(queue);
            CL_CHECK_ERROR (err);
            evKernel.FillTimingInfo();
            totalReduceTime += (evKernel.SubmitEndRuntime() / 1.e9);
        }

        err = clEnqueueReadBuffer(queue, d_odata, true, 0,
                num_blocks*sizeof(T), h_odata, 0, NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        double totalTransfer = (inputTransfer + evTransfer.StartEndRuntime()) /
                1.e9;
        
        T local_result = 0.0f, global_result = 0.0f;
        
        // Start a wallclock timer for MPI
        int TH_global = Timer::Start();
        
        // Perform reduction of block sums and MPI allreduce call
        for (int m = 0; m < iters; m++)
        {        
            local_result = 0.0f;

            for (int i=0; i<num_blocks; i++)
            {
                local_result += h_odata[i];
            }
            global_result = 0.0f;
            globalReduction(&local_result, &global_result);
        }
        double mpi_time = Timer::Stop(TH_global,"global all reduce") / iters;

        // Compute local reference solution
        T cpu_result = reduceCPU<T>(h_idata, size);
        // Use some error threshold for floating point rounding
        double threshold = 1.0e-6;
        T diff = fabs(local_result - cpu_result);

        if (diff > threshold)
        {
            cout << "Error in local reduction detected in rank " 
                 << mpi_rank << "\n";
            cout << "Diff: " << diff << endl;
        }
        
        if (global_result != (mpi_size * local_result))
        {
            cout << "Test Failed, error in global all reduce detected in rank "
                 << mpi_rank << endl;
        } 
        else 
        {
            if (mpi_rank == 0)
            {
                cout << "Test Passed.\n";
            }
        }
        // Calculate results
        char atts[1024];
        sprintf(atts, "%d_itemsPerRank",size);
        double local_gbytes = (double)(size*sizeof(T))/(1000.*1000.*1000.);
        double global_gbytes = local_gbytes * mpi_size;
        totalReduceTime /= iters; // use average time over the iterations
        resultDB.AddResult(testName+"-Kernel", atts, "GB/s", 
            global_gbytes / totalReduceTime);
        resultDB.AddResult(testName+"-Kernel+PCIe", atts, "GB/s",
            global_gbytes / (totalReduceTime + totalTransfer));
        resultDB.AddResult(testName+"-MPI_Allreduce",  atts, "GB/s",
            (sizeof(T)*mpi_size*1.e-9) / (mpi_time));
        resultDB.AddResult(testName+"-Overall", atts, "GB/s", 
            global_gbytes / (totalReduceTime + totalTransfer + mpi_time));
        }

    err = clEnqueueUnmapMemObject(queue, h_i, h_idata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_o, h_odata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_i);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_o);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_idata);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_odata);
    CL_CHECK_ERROR(err);
    err = clReleaseProgram(prog);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(reduce);
    CL_CHECK_ERROR(err);
}
예제 #20
0
void dump(OptionParser& op)
{	
    int i, j;
    void* work;
    T2* source, * result;
    unsigned long bytes = 0;

    if (op.getOptionInt("MB") == 0) {
        int probSizes[4] = { 1, 8, 96, 256 };
	int sizeIndex = op.getOptionInt("size")-1;
	if (sizeIndex < 0 || sizeIndex >= 4) {
	    cerr << "Invalid size index specified\n";
	    exit(-1);
	}
        bytes = probSizes[sizeIndex];
    } else {
        bytes = op.getOptionInt("MB");
    }

    // Convert to MB
    bytes *= 1024 * 1024;

    bool do_dp = dp<T2>();
    init(op, do_dp);

    // now determine how much available memory will be used
    int half_n_ffts = bytes / (512*sizeof(T2)*2);
    int n_ffts = half_n_ffts * 2;
    int half_n_cmplx = half_n_ffts * 512;
    unsigned long used_bytes = half_n_cmplx * 2 * sizeof(T2);
    double N = half_n_cmplx*2;
    
    fprintf(stderr, "used_bytes=%d, N=%g\n", used_bytes, N);

    // allocate host and device memory
    allocHostBuffer((void**)&source, used_bytes);
    allocHostBuffer((void**)&result, used_bytes);

    // init host memory...
    for (i = 0; i < half_n_cmplx; i++) {
        source[i].x = (rand()/(float)RAND_MAX)*2-1;
        source[i].y = (rand()/(float)RAND_MAX)*2-1;
        source[i+half_n_cmplx].x = source[i].x;
        source[i+half_n_cmplx].y = source[i].y;
    }

    // alloc device memory
    allocDeviceBuffer(&work, used_bytes);

    copyToDevice(work, source, used_bytes);

    fprintf(stdout, "INITIAL:\n");
    for (i = 0; i < N; i++) {
        fprintf(stdout, "(%g, %g)\n", source[i].x, source[i].y);
    }
    
    forward(work, n_ffts);
    copyFromDevice(result, work, used_bytes);
        
    fprintf(stdout, "FORWARD:\n");
    for (i = 0; i < N; i++) {
        fprintf(stdout, "(%g, %g)\n", result[i].x, result[i].y);
    }
    
    inverse(work, n_ffts);
    copyFromDevice(result, work, used_bytes);
        
    fprintf(stdout, "\nINVERSE:\n");
    for (i = 0; i < N; i++) {
        fprintf(stdout, "(%g, %g)\n", result[i].x, result[i].y);
    }

    freeDeviceBuffer(work);
    freeHostBuffer(source);
    freeHostBuffer(result);
}
예제 #21
0
파일: Spmv.cpp 프로젝트: bart-utahman/shoc
void csrTest(cl_device_id dev, cl_context ctx, string compileFlags,
             cl_command_queue queue, ResultDatabase& resultDB, OptionParser& op,
             floatType* h_val, int* h_cols, int* h_rowDelimiters,
             floatType* h_vec, floatType* h_out, int numRows, int numNonZeroes,
             floatType* refOut, bool padded, const size_t maxImgWidth)
{
    if (devSupportsImages)
    {
        char texflags[64];
        sprintf(texflags," -DUSE_TEXTURE -DMAX_IMG_WIDTH=%ld", maxImgWidth);
        compileFlags+=string(texflags);
    }
    // Set up OpenCL Program Object
    int err = 0;

    cl_program prog = clCreateProgramWithSource(ctx, 1, &cl_source_spmv, NULL,
            &err);
    CL_CHECK_ERROR(err);

    // Build the openCL kernels
    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    // CL_CHECK_ERROR(err);  // if we check and fail here, we never get to see
                            // the OpenCL compiler's build log

    // If there is a build error, print the output and return
    if (err != CL_SUCCESS)
    {
        char log[5000];
        size_t retsize = 0;
        err = clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 5000
                * sizeof(char), log, &retsize);
        CL_CHECK_ERROR(err);
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

      // Device data structures
      cl_mem d_val, d_vec, d_out;
      cl_mem d_cols, d_rowDelimiters;

      // Allocate device memory
      d_val = clCreateBuffer(ctx, CL_MEM_READ_WRITE, numNonZeroes *
          sizeof(clFloatType), NULL, &err);
      CL_CHECK_ERROR(err);
      d_cols = clCreateBuffer(ctx, CL_MEM_READ_WRITE, numNonZeroes *
          sizeof(cl_int), NULL, &err);
      CL_CHECK_ERROR(err);
      int imgHeight = 0;
      if (devSupportsImages)
      {
          imgHeight=(numRows+maxImgWidth-1)/maxImgWidth;
          cl_image_format fmt; fmt.image_channel_data_type=CL_FLOAT;
          if(sizeof(floatType)==4)
              fmt.image_channel_order=CL_R;
          else
              fmt.image_channel_order=CL_RG;
          d_vec = clCreateImage2D( ctx, CL_MEM_READ_ONLY, &fmt, maxImgWidth,
              imgHeight, 0, NULL, &err);
          CL_CHECK_ERROR(err);
      } else {
          d_vec = clCreateBuffer(ctx, CL_MEM_READ_WRITE, numRows *
              sizeof(clFloatType), NULL, &err);
          CL_CHECK_ERROR(err);
      }
      d_out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, numRows *
          sizeof(clFloatType), NULL, &err);
      CL_CHECK_ERROR(err);
      d_rowDelimiters = clCreateBuffer(ctx, CL_MEM_READ_WRITE, (numRows+1) *
          sizeof(cl_int), NULL, &err);
      CL_CHECK_ERROR(err);

      // Setup events for timing
      Event valTransfer("transfer Val data over PCIe bus");
      Event colsTransfer("transfer cols data over PCIe bus");
      Event vecTransfer("transfer vec data over PCIe bus");
      Event rowDelimitersTransfer("transfer rowDelimiters data over PCIe bus");

      // Transfer data to device
      err = clEnqueueWriteBuffer(queue, d_val, true, 0, numNonZeroes *
          sizeof(floatType), h_val, 0, NULL, &valTransfer.CLEvent());
      CL_CHECK_ERROR(err);
      err = clEnqueueWriteBuffer(queue, d_cols, true, 0, numNonZeroes *
          sizeof(int), h_cols, 0, NULL, &colsTransfer.CLEvent());
      CL_CHECK_ERROR(err);

      if (devSupportsImages)
      {
          size_t offset[3]={0};
          size_t size[3]={maxImgWidth,(size_t)imgHeight,1};
          err = clEnqueueWriteImage(queue,d_vec, true, offset, size,
              0, 0, h_vec, 0, NULL, &vecTransfer.CLEvent());
          CL_CHECK_ERROR(err);
      } else
      {
          err = clEnqueueWriteBuffer(queue, d_vec, true, 0, numRows *
              sizeof(floatType), h_vec, 0, NULL, &vecTransfer.CLEvent());
          CL_CHECK_ERROR(err);
      }

      err = clEnqueueWriteBuffer(queue, d_rowDelimiters, true, 0, (numRows+1) *
          sizeof(int), h_rowDelimiters, 0, NULL,
          &rowDelimitersTransfer.CLEvent());
      CL_CHECK_ERROR(err);
      err = clFinish(queue);
      CL_CHECK_ERROR(err);

      valTransfer.FillTimingInfo();
      colsTransfer.FillTimingInfo();
      vecTransfer.FillTimingInfo();
      rowDelimitersTransfer.FillTimingInfo();

      double iTransferTime = valTransfer.StartEndRuntime() +
                            colsTransfer.StartEndRuntime() +
                             vecTransfer.StartEndRuntime() +
                   rowDelimitersTransfer.StartEndRuntime();

      int passes = op.getOptionInt("passes");
      int iters  = op.getOptionInt("iterations");

      // Results description info
      char atts[TEMP_BUFFER_SIZE];
      sprintf(atts, "%d_elements_%d_rows", numNonZeroes, numRows);
      string prefix = "";
      prefix += (padded) ? "Padded_" : "";
      double gflop = 2 * (double) numNonZeroes;
      cout << "CSR Scalar Kernel\n";
      Event kernelExec("kernel Execution");

      // Set up CSR Kernels
      cl_kernel csrScalar, csrVector;
      csrScalar  = clCreateKernel(prog, "spmv_csr_scalar_kernel", &err);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 0, sizeof(cl_mem), (void*) &d_val);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 1, sizeof(cl_mem), (void*) &d_vec);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 2, sizeof(cl_mem), (void*) &d_cols);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 3, sizeof(cl_mem),
          (void*) &d_rowDelimiters);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 4, sizeof(cl_int), (void*) &numRows);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrScalar, 5, sizeof(cl_mem), (void*) &d_out);
      CL_CHECK_ERROR(err);

      csrVector = clCreateKernel(prog, "spmv_csr_vector_kernel", &err);

      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 0, sizeof(cl_mem), (void*) &d_val);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 1, sizeof(cl_mem), (void*) &d_vec);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 2, sizeof(cl_mem), (void*) &d_cols);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 3, sizeof(cl_mem),
          (void*) &d_rowDelimiters);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 4, sizeof(cl_int), (void*) &numRows);
      CL_CHECK_ERROR(err);
      err = clSetKernelArg(csrVector, 5, sizeof(cl_mem), (void*) &d_out);
      CL_CHECK_ERROR(err);

      // Append correct suffix to resultsDB entry
      string suffix;
      if (sizeof(floatType) == sizeof(float))
      {
          suffix = "-SP";
      }
      else
      {
          suffix = "-DP";
      }

      const size_t scalarGlobalWSize = numRows;
      size_t localWorkSize = BLOCK_SIZE;

      for (int k = 0; k < passes; k++)
      {
          double scalarKernelTime = 0.0;
          // Run Scalar Kernel
          for (int j = 0; j < iters; j++)
          {
              err = clEnqueueNDRangeKernel(queue, csrScalar, 1, NULL,
                   &scalarGlobalWSize, &localWorkSize, 0, NULL,
                   &kernelExec.CLEvent());
              CL_CHECK_ERROR(err);
              err = clFinish(queue);
              CL_CHECK_ERROR(err);
              kernelExec.FillTimingInfo();
              scalarKernelTime += kernelExec.StartEndRuntime();
          }

          // Transfer data back to host
          Event outTransfer("d->h data transfer");
          err = clEnqueueReadBuffer(queue, d_out, true, 0, numRows *
              sizeof(floatType), h_out, 0, NULL, &outTransfer.CLEvent());
          CL_CHECK_ERROR(err);
          err = clFinish(queue);
          CL_CHECK_ERROR(err);
          outTransfer.FillTimingInfo();
          double oTransferTime = outTransfer.StartEndRuntime();

          // Compare reference solution to GPU result
          if (! verifyResults(refOut, h_out, numRows, k))
          {
               return;  // If results don't match, don't report performance
          }
          scalarKernelTime = scalarKernelTime / (double)iters;
          string testName = prefix+"CSR-Scalar"+suffix;
          double totalTransfer = iTransferTime + oTransferTime;
          resultDB.AddResult(testName, atts, "Gflop/s",
              gflop/(scalarKernelTime));
          resultDB.AddResult(testName+"_PCIe", atts, "Gflop/s",
              gflop / (scalarKernelTime+totalTransfer));
      }

      // Clobber correct answer, so we can be sure the vector kernel is correct
      err = clEnqueueWriteBuffer(queue, d_out, true, 0, numRows *
          sizeof(floatType), h_vec, 0, NULL, NULL);
      CL_CHECK_ERROR(err);

      cout << "CSR Vector Kernel\n";
      // Verify Local work group size
      size_t maxLocal = getMaxWorkGroupSize(ctx, csrVector);
      if (maxLocal < 32)
      {
         cout << "Warning: CSRVector requires a work group size >= 32" << endl;
         cout << "Skipping this kernel." << endl;
         err = clReleaseMemObject(d_rowDelimiters);
         CL_CHECK_ERROR(err);
         err = clReleaseMemObject(d_vec);
         CL_CHECK_ERROR(err);
         err = clReleaseMemObject(d_out);
         CL_CHECK_ERROR(err);
         err = clReleaseMemObject(d_val);
         CL_CHECK_ERROR(err);
         err = clReleaseMemObject(d_cols);
         CL_CHECK_ERROR(err);
         err = clReleaseKernel(csrScalar);
         CL_CHECK_ERROR(err);
         err = clReleaseKernel(csrVector);
         CL_CHECK_ERROR(err);
         err = clReleaseProgram(prog);
         CL_CHECK_ERROR(err);
         return;
      }
      localWorkSize = VECTOR_SIZE;
      while (localWorkSize+VECTOR_SIZE <= maxLocal &&
          localWorkSize+VECTOR_SIZE <= BLOCK_SIZE)
      {
         localWorkSize += VECTOR_SIZE;
      }
      const size_t vectorGlobalWSize = numRows * VECTOR_SIZE; // 1 warp per row

      for (int k = 0; k < passes; k++)
      {
          // Run Vector Kernel
          double vectorKernelTime = 0.0;
          for (int j = 0; j < iters; j++)
          {
             err = clEnqueueNDRangeKernel(queue, csrVector, 1, NULL,
                  &vectorGlobalWSize, &localWorkSize, 0, NULL,
                  &kernelExec.CLEvent());
             CL_CHECK_ERROR(err);
             err = clFinish(queue);
             CL_CHECK_ERROR(err);
             kernelExec.FillTimingInfo();
             vectorKernelTime += kernelExec.StartEndRuntime();
          }

         Event outTransfer("d->h data transfer");
         err = clEnqueueReadBuffer(queue, d_out, true, 0, numRows *
             sizeof(floatType), h_out, 0, NULL, &outTransfer.CLEvent());
         CL_CHECK_ERROR(err);
         err = clFinish(queue);
         CL_CHECK_ERROR(err);
         outTransfer.FillTimingInfo();
         double oTransferTime = outTransfer.StartEndRuntime();

          // Compare reference solution to GPU result
          if (! verifyResults(refOut, h_out, numRows, k))
          {
              return;  // If results don't match, don't report performance
          }
          vectorKernelTime = vectorKernelTime / (double)iters;
          string testName = prefix+"CSR-Vector"+suffix;
          double totalTransfer = iTransferTime + oTransferTime;
          resultDB.AddResult(testName, atts, "Gflop/s", gflop/vectorKernelTime);
          resultDB.AddResult(testName+"_PCIe", atts, "Gflop/s",
                            gflop/(vectorKernelTime+totalTransfer));
      }

      // Free device memory
      err = clReleaseMemObject(d_rowDelimiters);
      CL_CHECK_ERROR(err);
      err = clReleaseMemObject(d_vec);
      CL_CHECK_ERROR(err);
      err = clReleaseMemObject(d_out);
      CL_CHECK_ERROR(err);
      err = clReleaseMemObject(d_val);
      CL_CHECK_ERROR(err);
      err = clReleaseMemObject(d_cols);
      CL_CHECK_ERROR(err);
      err = clReleaseKernel(csrScalar);
      CL_CHECK_ERROR(err);
      err = clReleaseKernel(csrVector);
      CL_CHECK_ERROR(err);
      err = clReleaseProgram(prog);
      CL_CHECK_ERROR(err);
}
예제 #22
0
파일: main.cpp 프로젝트: ManavA/shoc
// ****************************************************************************
// Method:  main()
//
// Purpose:
//   serial and parallel main for OpenCL level0 benchmarks
//
// Arguments:
//   argc, argv
//
// Programmer:  SHOC Team
// Creation:    The Epoch
//
// Modifications:
//   Jeremy Meredith, Tue Jan 12 15:09:33 EST 2010
//   Changed the way device selection works.  It now defaults to the device
//   index corresponding to the process's rank within a node if no devices
//   are specified on the command command line, and otherwise, round-robins
//   the list of devices among the tasks.
//
//   Gabriel Marin, Tue Jun 01 15:38 EST 2010
//   Check that we have valid (not NULL) context and queue objects before
//   running the benchmarks. Errors inside CreateContextFromSingleDevice or
//   CreateCommandQueueForContextAndDevice were not propagated out to the main
//   program.
//
//   Jeremy Meredith, Wed Nov 10 14:20:47 EST 2010
//   Split timing reports into detailed and summary.  For serial code, we
//   report all trial values, and for parallel, skip the per-process vals.
//   Also detect and print outliers from parallel runs.
//
// ****************************************************************************
int main(int argc, char *argv[])
{
    int ret = 0;

    try
    {
#ifdef PARALLEL
        int rank, size;
        MPI_Init(&argc,&argv);
        MPI_Comm_size(MPI_COMM_WORLD, &size);
        MPI_Comm_rank(MPI_COMM_WORLD, &rank);
        cout << "MPI Task "<< rank << "/" << size - 1 << " starting....\n";
#endif

        OptionParser op;

        //Add shared options to the parser
        op.addOption("platform", OPT_INT, "0", "specify OpenCL platform to use",
                'p');
        op.addOption("device", OPT_VECINT, "", "specify device(s) to run on", 'd');
        op.addOption("passes", OPT_INT, "10", "specify number of passes", 'n');
        op.addOption("size", OPT_VECINT, "1", "specify problem size", 's');
        op.addOption("infoDevices", OPT_BOOL, "",
                "show info for available platforms and devices", 'i');
        op.addOption("verbose", OPT_BOOL, "", "enable verbose output", 'v');
        op.addOption("quiet", OPT_BOOL, "", "write minimum necessary to standard output", 'q');

        addBenchmarkSpecOptions(op);

        if (!op.parse(argc, argv))
        {
#ifdef PARALLEL
            if (rank == 0)
                op.usage();
            MPI_Finalize();
#else
            op.usage();
#endif
            return (op.HelpRequested() ? 0 : 1 );
        }

        if (op.getOptionBool("infoDevices"))
        {
#define DEBUG_DEVICE_CONTAINER 0
#ifdef PARALLEL
            // execute following code only if I am the process of lowest
            // rank on this node
            NodeInfo NI;
            int mynoderank = NI.nodeRank();
            if (mynoderank==0)
            {
                int nlrrank, nlrsize;
                MPI_Comm nlrcomm = NI.getNLRComm();
                MPI_Comm_size(nlrcomm, &nlrsize);
                MPI_Comm_rank(nlrcomm, &nlrrank);

                OpenCLNodePlatformContainer ndc1;
                OpenCLMultiNodeContainer localMnc(ndc1);
                localMnc.doMerge (nlrrank, nlrsize, nlrcomm);
                if (rank==0)  // I am the global rank 0, print all configurations
                    localMnc.Print (cout);
            }
#else
            OpenCLNodePlatformContainer ndc1;
            ndc1.Print (cout);
#if DEBUG_DEVICE_CONTAINER
            OpenCLMultiNodeContainer mnc1(ndc1), mnc2;
            mnc1.Print (cout);
            ostringstream oss;
            mnc1.writeObject (oss);
            std::string temp(oss.str());
            cout << "Serialized MultiNodeContainer:\n" << temp;
            istringstream iss(temp);
            mnc2.readObject (iss);
            cout << "Unserialized object2:\n";
            mnc2.Print (cout);
            mnc1.merge (mnc2);
            cout << "==============\nObject1 after merging 1:\n";
            mnc1.Print (cout);
            mnc1.merge (mnc2);
            cout << "==============\nObject1 after merging 2:\n";
            mnc1.Print (cout);
#endif  // DEBUG
#endif  // PARALLEL
            return (0);
        }

        bool verbose = op.getOptionBool("verbose");

        // The device option supports specifying more than one device
        // for now, just choose the first one.
        int platform = op.getOptionInt("platform");

#ifdef PARALLEL
        NodeInfo ni;
        int myNodeRank = ni.nodeRank();
        if (verbose)
        cout << "Global rank "<<rank<<" is local rank "<<myNodeRank << endl;
#else
        int myNodeRank = 0;
#endif

        // If they haven't specified any devices, assume they
        // want the process with in-node rank N to use device N
        int deviceIdx = myNodeRank;

        // If they have, then round-robin the list of devices
        // among the processes on a node.
        vector<long long> deviceVec = op.getOptionVecInt("device");
        if (deviceVec.size() > 0)
        {
        int len = deviceVec.size();
            deviceIdx = deviceVec[myNodeRank % len];
        }

        // Check for an erroneous device
        if (deviceIdx >= GetNumOclDevices(platform)) {
            cerr << "Warning: device index: " << deviceIdx
                 << " out of range, defaulting to device 0.\n";
            deviceIdx = 0;
        }

        // Initialization
        if (verbose) cout << ">> initializing\n";
        cl_device_id devID = ListDevicesAndGetDevice(platform, deviceIdx);
        cl_int clErr;
        cl_context ctx = clCreateContext( NULL,     // properties
                                            1,      // number of devices
                                            &devID, // device
                                            NULL,   // notification function
                                            NULL,
                                            &clErr );
        CL_CHECK_ERROR(clErr);
        cl_command_queue queue = clCreateCommandQueue( ctx,
                                                        devID,
                                                        CL_QUEUE_PROFILING_ENABLE,
                                                        &clErr );
        CL_CHECK_ERROR(clErr);
        ResultDatabase resultDB;

        // Run the benchmark
        RunBenchmark(devID, ctx, queue, resultDB, op);

        clReleaseCommandQueue( queue );
        clReleaseContext( ctx );


#ifndef PARALLEL
        resultDB.DumpDetailed(cout);
#else
        ParallelResultDatabase pardb;
        pardb.MergeSerialDatabases(resultDB,MPI_COMM_WORLD);
        if (rank==0)
        {
            pardb.DumpSummary(cout);
            pardb.DumpOutliers(cout);
        }
#endif
    }
    catch( std::exception& e )
    {
        std::cerr << e.what() << std::endl;
        ret = 1;
    }
    catch( ... )
    {
        std::cerr << "unrecognized exception caught" << std::endl;
        ret = 1;
    }

#ifdef PARALLEL
    MPI_Finalize();
#endif

    return ret;
}
예제 #23
0
파일: Sort.cpp 프로젝트: ellen-hl/shoc-mic
void RunTest(string testName, ResultDatabase &resultDB, OptionParser &op)
{
    int probSizes[4] = { 1, 8, 48, 96 };

    int size = probSizes[op.getOptionInt("size")-1];
    
    // Convert to MiB
    size = (size*1024*1024)/sizeof(T);
    
    // Create input data on CPU
    unsigned int bytes = size * sizeof(T);

    // Allocate Host Memory
    __declspec(target(MIC)) static T *hkey, *outkey;
    __declspec(target(MIC)) static T *hvalue, *outvalue;

    hkey   = (T*)_mm_malloc(bytes,ALIGN);
    hvalue = (T*)_mm_malloc(bytes,ALIGN);

    outkey   = (T*)_mm_malloc(bytes,ALIGN);
    outvalue = (T*)_mm_malloc(bytes,ALIGN);


    // Initialize host memory
    cout << "Initializing host memory." << endl;

    srand(time(NULL));
    for (int i = 0; i < size; i++)
    {
        hkey[i] = hvalue[i]= (i+255) % 1089; // Fill with some pattern
    }

    int micdev = op.getOptionInt("target");
    int iters = op.getOptionInt("passes");
    int numThreads = op.getOptionInt("nthreads");

    cout << "nthreads   = " <<numThreads<< endl;

    cout << "Running benchmark" << endl;
    for(int it=0;it<iters;it++)
    {

        // Allocating buffer on card
        #pragma offload target(mic:micdev) in(hkey:length(size)  free_if(0)) \
                in(hvalue:length(size) free_if(0))\
                out(outkey:length(size) free_if(0))\
                out(outvalue:length(size) free_if(0))
        {
        }

        double start = curr_second();
        // Get data transfer time
        #pragma offload target(mic:micdev) in(hkey:length(size) alloc_if(0) \
                free_if(0)) in(hvalue:length(size) alloc_if(0) free_if(0))
        {
        }

        float transferTime = curr_second()-start;
        double totalRunTime = 0.0f;
        start = curr_second();
        #pragma offload target(mic:micdev) nocopy(hkey:length(size) \
                alloc_if(0) free_if(0))                             \
                nocopy(hvalue:length(size) alloc_if(0) free_if(0))  \
                nocopy(outkey:length(size) alloc_if(0) free_if(0))  \
                nocopy(outvalue:length(size) alloc_if(0) free_if(0))\
                in(numThreads)
        {
            sortKernel<T>(hkey, hvalue, outkey, outvalue, size, numThreads);
        }
        totalRunTime = curr_second()-start;

        #pragma offload target(mic:micdev) nocopy(hkey:length(size) \
                alloc_if(0) free_if(1))                             \
                nocopy(hvalue:length(size) alloc_if(0) free_if(1))  \
                out(outkey:length(size) alloc_if(0))                \
                out(outvalue:length(size) alloc_if(0))
        {
        }

        // If results aren't correct, don't report perf numbers
        if (!verifyResult<T>(outkey, outvalue, size))
        {
            return;
        }

        char atts[1024];
        double avgTime = (totalRunTime / (double) iters);
        sprintf(atts, "%d items", size);
        double gb = (double)(size * sizeof(T)) / (1000. * 1000. * 1000.);
        resultDB.AddResult(testName, atts, "GB/s", gb / avgTime);
        resultDB.AddResult(testName+"_PCIe", atts, "GB/s",
                gb / (avgTime + transferTime));
        resultDB.AddResult(testName+"_Parity", atts, "N",
                transferTime / avgTime);

    }
    // Clean up
    _mm_free(hkey);
    _mm_free(hvalue);

}
예제 #24
0
파일: S3D.cpp 프로젝트: dylanzika/shoc
void RunTest(const string& testName, cl_device_id dev, cl_context ctx,
             cl_command_queue queue, ResultDatabase &resultDB,
             OptionParser &op, string& compileFlags)
{
    int n_species = 22;
    int i, j, err;

    int probSizes_SP[4] = { 24, 32, 40, 48};
    int probSizes_DP[4] = { 16, 24, 32, 40};
    int *probSizes = (sizeof(T) == sizeof(double)) ? probSizes_DP : probSizes_SP;
    int size = probSizes[op.getOptionInt("size")-1];

    // The number of grid points
    int n = size * size * size;

    // For now these conversion factors are just 1
    T pconv = 1.0;    // 1418365.88544;
    T tconv = 1.0;    //120.0;
    T rateconv = 1.0; //11.0393507649917;

    // Host copies of data
    T* h_t     = new T[n];
    T* h_p     = new T[n];
    T* h_y     = new T[n*n_species];
    T* h_wdot  = new T[n*n_species];
    T* h_molwt = new T[n_species];

    // Device data
    cl_mem d_t;    // Temperatures array
    cl_mem d_p;    // Pressures array
    cl_mem d_y;    // Input variables
    cl_mem d_wdot; // Output variables

    // intermediate variables
    cl_mem d_rf, d_rb, d_rklow, d_c, d_a, d_eg, d_molwt;

    // Initialize host memory
    for (i=0; i<n; i++)
    {
        h_p[i] = 1.0132e6;
        h_t[i] = 1000.0;
    }

    for (j=0; j<22; j++)
    {
        for (i=0; i<n; i++)
        {
            h_y[(j*n)+i]= 0.0;
            if (j==14)
                h_y[(j*n)+i] = 0.064;
            if (j==3)
                h_y[(j*n)+i] = 0.218;
            if (j==21)
                h_y[(j*n)+i] = 0.718;
        }
    }

    for (int i=0; i<n_species; i++)
    {
        h_molwt[i] = 1.0f;
    }
//    // Initialize molecular weights
//    h_molwt[0]= 2.01594E-03;
//    h_molwt[1]= 1.00797E-03;
//    h_molwt[2]= 1.59994E-02;
//    h_molwt[3]= 3.19988E-02;
//    h_molwt[4]= 1.700737E-02;
//    h_molwt[5]= 1.801534E-02;
//    h_molwt[6]= 3.300677E-02;
//    h_molwt[7]= 3.401473999999999E-02;
//    h_molwt[8]= 1.503506E-02;
//    h_molwt[9] = 1.604303E-02;
//    h_molwt[10] = 2.801055E-02;
//    h_molwt[11] = 4.400995E-02;
//    h_molwt[12] = 3.002649E-02;
//    h_molwt[13] = 2.603824E-02;
//    h_molwt[14] = 2.805418E-02;
//    h_molwt[15] = 3.007012E-02;
//    h_molwt[16] = 4.102967E-02;
//    h_molwt[17] = 4.203764E-02;
//    h_molwt[18] = 4.405358E-02;
//    h_molwt[19] = 4.10733E-02;
//    h_molwt[20] = 4.208127E-02;
//    h_molwt[21] = 2.80134E-02;

    // Allocate device memory
    size_t base = n * sizeof(T);
    clMalloc(d_t, base);
    clMalloc(d_p, base);
    clMalloc(d_y, n_species*base);
    clMalloc(d_wdot, n_species*base);
    clMalloc(d_rf, 206*base);
    clMalloc(d_rb, 206*base);
    clMalloc(d_rklow, 21*base);
    clMalloc(d_c, C_SIZE*base);
    clMalloc(d_a, A_SIZE*base);
    clMalloc(d_eg, EG_SIZE*base);
    clMalloc(d_molwt, n_species*sizeof(T));

    // Copy over input params
    long inputTransferTime = 0;
    Event evTransfer("PCIe Transfer");

    clMemtoDevice(d_t, h_t, base);
    evTransfer.FillTimingInfo();
    inputTransferTime += evTransfer.StartEndRuntime();

    clMemtoDevice(d_p, h_p, base);
    evTransfer.FillTimingInfo();
    inputTransferTime += evTransfer.StartEndRuntime();

    clMemtoDevice(d_y, h_y, n_species*base);
    evTransfer.FillTimingInfo();
    inputTransferTime += evTransfer.StartEndRuntime();

    clMemtoDevice(d_molwt, h_molwt, n_species*sizeof(T));
    evTransfer.FillTimingInfo();
    inputTransferTime += evTransfer.StartEndRuntime();

    // Set up macros
    compileFlags += "-DDIM="  + toString(size) + " " +
                    "-DN_GP=" + toString(n)    + " ";

    unsigned int passes = op.getOptionInt("passes");
    for (unsigned int i = 0; i < passes; i++)
    {
        size_t globalWorkSize = n;
        size_t localWorkSize = 128;

        // -------------------- phase 1 -----------------

        // Setup Program Objects (phase 1)
        clProg(gr_prog, cl_source_gr_base);
        clProg(rdsmh_prog, cl_source_rdsmh);
        clProg(ratt_prog, cl_source_ratt);
        clProg(ratt2_prog, cl_source_ratt2);
        clProg(ratt3_prog, cl_source_ratt3);
        clProg(ratt4_prog, cl_source_ratt4);
        clProg(ratt5_prog, cl_source_ratt5);
        clProg(ratt6_prog, cl_source_ratt6);
        clProg(ratt7_prog, cl_source_ratt7);
        clProg(ratt8_prog, cl_source_ratt8);
        clProg(ratt9_prog, cl_source_ratt9);
        clProg(ratt10_prog, cl_source_ratt10);
        clProg(ratx_prog, cl_source_ratx);
        clProg(ratxb_prog, cl_source_ratxb);
        clProg(ratx2_prog, cl_source_ratx2);
        clProg(ratx4_prog, cl_source_ratx4);

        // Build the kernels (phase 1)
        cout << "Compiling kernels (phase 1)...";
        cout.flush();

        clBuild(gr_prog);
        clBuild(rdsmh_prog);
        clBuild(ratt_prog);
        clBuild(ratt2_prog);
        clBuild(ratt3_prog);
        clBuild(ratt4_prog);
        clBuild(ratt5_prog);
        clBuild(ratt6_prog);
        clBuild(ratt7_prog);
        clBuild(ratt8_prog);
        clBuild(ratt9_prog);
        clBuild(ratt10_prog);
        clBuild(ratx_prog);
        clBuild(ratxb_prog);
        clBuild(ratx2_prog);
        clBuild(ratx4_prog);

        cout << "done." << endl;

        // Extract out kernel objects (phase 1)
        cout << "Generating OpenCL Kernel Objects (phase 1)...";
        cout.flush();

        // GR Base Kernels
        cl_kernel grBase_kernel = clCreateKernel(gr_prog, "gr_base", &err);
        CL_CHECK_ERROR(err);

        // RDSMH Kernels
        cl_kernel rdsmh_kernel = clCreateKernel(rdsmh_prog, "rdsmh_kernel",
                &err);
        CL_CHECK_ERROR(err);

        // RATT Kernels
        cl_kernel ratt_kernel = clCreateKernel(ratt_prog, "ratt_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt2_kernel = clCreateKernel(ratt2_prog, "ratt2_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt3_kernel = clCreateKernel(ratt3_prog, "ratt3_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt4_kernel = clCreateKernel(ratt4_prog, "ratt4_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt5_kernel = clCreateKernel(ratt5_prog, "ratt5_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt6_kernel = clCreateKernel(ratt6_prog, "ratt6_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt7_kernel = clCreateKernel(ratt7_prog, "ratt7_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt8_kernel = clCreateKernel(ratt8_prog, "ratt8_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt9_kernel = clCreateKernel(ratt9_prog, "ratt9_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratt10_kernel = clCreateKernel(ratt10_prog, "ratt10_kernel",
                &err);
        CL_CHECK_ERROR(err);

        // RATX Kernels
        cl_kernel ratx_kernel = clCreateKernel(ratx_prog, "ratx_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratxb_kernel = clCreateKernel(ratxb_prog, "ratxb_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratx2_kernel = clCreateKernel(ratx2_prog, "ratx2_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel ratx4_kernel = clCreateKernel(ratx4_prog, "ratx4_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cout << "done." << endl;

        //Set kernel arguments (phase 1)
        err = clSetKernelArg(grBase_kernel, 0, sizeof(cl_mem), (void*)&d_p);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(grBase_kernel, 1, sizeof(cl_mem), (void*)&d_t);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(grBase_kernel, 2, sizeof(cl_mem), (void*)&d_y);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(grBase_kernel, 3, sizeof(cl_mem), (void*)&d_c);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(grBase_kernel, 4, sizeof(T), (void*)&tconv);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(grBase_kernel, 5, sizeof(T), (void*)&pconv);
        CL_CHECK_ERROR(err);

        err = clSetKernelArg(rdsmh_kernel, 0, sizeof(cl_mem), (void*)&d_t);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(rdsmh_kernel, 1, sizeof(cl_mem), (void*)&d_eg);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(rdsmh_kernel, 2, sizeof(T), (void*)&tconv);
        CL_CHECK_ERROR(err);

        err = clSetKernelArg(ratt_kernel, 0, sizeof(cl_mem), (void*)&d_t);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratt_kernel, 1, sizeof(cl_mem), (void*)&d_rf);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratt_kernel, 2, sizeof(T), (void*)&tconv);
        CL_CHECK_ERROR(err);
        clSetRattArg(ratt2_kernel);
        clSetRattArg(ratt3_kernel);
        clSetRattArg(ratt4_kernel);
        clSetRattArg(ratt5_kernel);
        clSetRattArg(ratt6_kernel);
        clSetRattArg(ratt7_kernel);
        clSetRattArg(ratt8_kernel);
        clSetRattArg(ratt9_kernel);
        err = clSetKernelArg(ratt10_kernel, 0, sizeof(cl_mem), (void*)&d_t);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratt10_kernel, 1, sizeof(cl_mem), (void*)&d_rklow);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratt10_kernel, 2, sizeof(T), (void*)&tconv);
        CL_CHECK_ERROR(err);

        clSetRatxArg(ratx_kernel);
        clSetRatxArg(ratxb_kernel);

        err = clSetKernelArg(ratx2_kernel, 0, sizeof(cl_mem), (void*)&d_c);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratx2_kernel, 1, sizeof(cl_mem), (void*)&d_rf);
        CL_CHECK_ERROR(err);

        err = clSetKernelArg(ratx4_kernel, 0, sizeof(cl_mem), (void*)&d_c);
        CL_CHECK_ERROR(err);
        err = clSetKernelArg(ratx4_kernel, 1, sizeof(cl_mem), (void*)&d_rb);
        CL_CHECK_ERROR(err);

        // Execute kernels (phase 1)
        cout << "Executing kernels (phase 1)...";
        cout.flush();

        Event evFirst_1("first kernel phase 1");
        Event evLast_1("last kernel phase 1");

        clLaunchKernelEv(grBase_kernel, evFirst_1.CLEvent());
        clLaunchKernel(ratt_kernel);
        clLaunchKernel(rdsmh_kernel);

        clLaunchKernel(ratt2_kernel);
        clLaunchKernel(ratt3_kernel);
        clLaunchKernel(ratt4_kernel);
        clLaunchKernel(ratt5_kernel);
        clLaunchKernel(ratt6_kernel);
        clLaunchKernel(ratt7_kernel);
        clLaunchKernel(ratt8_kernel);
        clLaunchKernel(ratt9_kernel);
        clLaunchKernel(ratt10_kernel);

        clLaunchKernel(ratx_kernel);
        clLaunchKernel(ratxb_kernel);
        clLaunchKernel(ratx2_kernel);
        clLaunchKernelEv(ratx4_kernel, evLast_1.CLEvent());

        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        cout << "done. " << endl;

        evFirst_1.FillTimingInfo();
        evLast_1.FillTimingInfo();
        double total_phase1 = evLast_1.EndTime() - evFirst_1.StartTime();

        // Release Kernels (phase 1)
        clReleaseKernel(grBase_kernel);
        clReleaseKernel(rdsmh_kernel);
        clReleaseKernel(ratt_kernel);
        clReleaseKernel(ratt2_kernel);
        clReleaseKernel(ratt3_kernel);
        clReleaseKernel(ratt4_kernel);
        clReleaseKernel(ratt5_kernel);
        clReleaseKernel(ratt6_kernel);
        clReleaseKernel(ratt7_kernel);
        clReleaseKernel(ratt8_kernel);
        clReleaseKernel(ratt9_kernel);
        clReleaseKernel(ratt10_kernel);
        clReleaseKernel(ratx_kernel);
        clReleaseKernel(ratxb_kernel);
        clReleaseKernel(ratx2_kernel);
        clReleaseKernel(ratx4_kernel);

        // Release Programs (phase 1)
        clReleaseProgram(gr_prog);
        clReleaseProgram(rdsmh_prog);
        clReleaseProgram(ratt_prog);
        clReleaseProgram(ratt2_prog);
        clReleaseProgram(ratt3_prog);
        clReleaseProgram(ratt4_prog);
        clReleaseProgram(ratt5_prog);
        clReleaseProgram(ratt6_prog);
        clReleaseProgram(ratt7_prog);
        clReleaseProgram(ratt8_prog);
        clReleaseProgram(ratt9_prog);
        clReleaseProgram(ratt10_prog);
        clReleaseProgram(ratx_prog);
        clReleaseProgram(ratxb_prog);
        clReleaseProgram(ratx2_prog);
        clReleaseProgram(ratx4_prog);

        // -------------------- phase 2 -----------------
        // Setup Program Objects (phase 2)
        clProg(qssa_prog, cl_source_qssa);
        clProg(qssab_prog, cl_source_qssab);
        clProg(qssa2_prog, cl_source_qssa2);
        clProg(rdwdot_prog, cl_source_rdwdot);
        clProg(rdwdot2_prog, cl_source_rdwdot2);
        clProg(rdwdot3_prog, cl_source_rdwdot3);
        clProg(rdwdot6_prog, cl_source_rdwdot6);
        clProg(rdwdot7_prog, cl_source_rdwdot7);
        clProg(rdwdot8_prog, cl_source_rdwdot8);
        clProg(rdwdot9_prog, cl_source_rdwdot9);
        clProg(rdwdot10_prog, cl_source_rdwdot10);

        // Build the kernels (phase 2)
        cout << "Compiling kernels (phase 2)...";
        cout.flush();

        clBuild(qssa_prog);
        clBuild(qssab_prog);
        clBuild(qssa2_prog);
        clBuild(rdwdot_prog);
        clBuild(rdwdot2_prog);
        clBuild(rdwdot3_prog);
        clBuild(rdwdot6_prog);
        clBuild(rdwdot7_prog);
        clBuild(rdwdot8_prog);
        clBuild(rdwdot9_prog);
        clBuild(rdwdot10_prog);

        cout << "done." << endl;

        // Extract out kernel objects (phase 2)
        cout << "Generating OpenCL Kernel Objects (phase 2)...";
        cout.flush();

        // QSSA Kernels
        cl_kernel qssa_kernel = clCreateKernel(qssa_prog, "qssa_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel qssab_kernel = clCreateKernel(qssab_prog, "qssab_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel qssa2_kernel = clCreateKernel(qssa2_prog, "qssa2_kernel",
                &err);
        CL_CHECK_ERROR(err);

        // RDWDOT Kernels
        cl_kernel rdwdot_kernel = clCreateKernel(rdwdot_prog, "rdwdot_kernel",
                &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot2_kernel = clCreateKernel(rdwdot2_prog,
                "rdwdot2_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot3_kernel = clCreateKernel(rdwdot3_prog,
                "rdwdot3_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot6_kernel = clCreateKernel(rdwdot6_prog,
                "rdwdot6_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot7_kernel = clCreateKernel(rdwdot7_prog,
                "rdwdot7_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot8_kernel = clCreateKernel(rdwdot8_prog,
                "rdwdot8_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot9_kernel = clCreateKernel(rdwdot9_prog,
                "rdwdot9_kernel", &err);
        CL_CHECK_ERROR(err);
        cl_kernel rdwdot10_kernel = clCreateKernel(rdwdot10_prog,
                "rdwdot10_kernel", &err);
        CL_CHECK_ERROR(err);
        cout << "done." << endl;

        //Set kernel arguments (phase 2)
        clSetQssaArg(qssa_kernel);
        clSetQssaArg(qssab_kernel);
        clSetQssaArg(qssa2_kernel);

        clSetRdwdotArg(rdwdot_kernel);
        clSetRdwdotArg(rdwdot2_kernel);
        clSetRdwdotArg(rdwdot3_kernel);
        clSetRdwdotArg(rdwdot6_kernel);
        clSetRdwdotArg(rdwdot7_kernel);
        clSetRdwdotArg(rdwdot8_kernel);
        clSetRdwdotArg(rdwdot9_kernel);
        clSetRdwdotArg(rdwdot10_kernel);

        // Execute kernels (phase 2)
        cout << "Executing kernels (phase 2)...";
        cout.flush();

        Event evFirst_2("first kernel phase 2");
        Event evLast_2("last kernel phase 2");

        clLaunchKernelEv(qssa_kernel, evFirst_2.CLEvent());
        clLaunchKernel(qssab_kernel);
        clLaunchKernel(qssa2_kernel);

        clLaunchKernel(rdwdot_kernel);
        clLaunchKernel(rdwdot2_kernel);
        clLaunchKernel(rdwdot3_kernel);
        clLaunchKernel(rdwdot6_kernel);
        clLaunchKernel(rdwdot7_kernel);
        clLaunchKernel(rdwdot8_kernel);
        clLaunchKernel(rdwdot9_kernel);
        clLaunchKernelEv(rdwdot10_kernel, evLast_2.CLEvent());

        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        cout << "done. " << endl;

        evFirst_2.FillTimingInfo();
        evLast_2.FillTimingInfo();
        double total_phase2 = evLast_2.EndTime() - evFirst_2.StartTime();

        // Release Kernels (phase 2)
        clReleaseKernel(qssa_kernel);
        clReleaseKernel(qssab_kernel);
        clReleaseKernel(qssa2_kernel);
        clReleaseKernel(rdwdot_kernel);
        clReleaseKernel(rdwdot2_kernel);
        clReleaseKernel(rdwdot3_kernel);
        clReleaseKernel(rdwdot6_kernel);
        clReleaseKernel(rdwdot7_kernel);
        clReleaseKernel(rdwdot8_kernel);
        clReleaseKernel(rdwdot9_kernel);
        clReleaseKernel(rdwdot10_kernel);

        // Release Programs (phase 2)
        clReleaseProgram(qssa_prog);
        clReleaseProgram(qssab_prog);
        clReleaseProgram(qssa2_prog);
        clReleaseProgram(rdwdot_prog);
        clReleaseProgram(rdwdot2_prog);
        clReleaseProgram(rdwdot3_prog);
        clReleaseProgram(rdwdot6_prog);
        clReleaseProgram(rdwdot7_prog);
        clReleaseProgram(rdwdot8_prog);
        clReleaseProgram(rdwdot9_prog);
        clReleaseProgram(rdwdot10_prog);

        // -------------------- timings -----------------

        double total = total_phase1 + total_phase2;
        // Estimate GFLOPs (roughly 10k flops / point)
        double gflops = (n*10000.) / total;
        // Copy results back
        err = clEnqueueReadBuffer(queue, d_wdot, true, 0,
                n*n_species*sizeof(T), h_wdot,
                0, NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        double totalTransferTime = inputTransferTime +
            evTransfer.StartEndRuntime();
        double gflops_pcie = (n*10000.) / (total + totalTransferTime);

        resultDB.AddResult(testName, "cubic", "GFLOPS", gflops);
        resultDB.AddResult(testName+"_PCIe", "cubic", "GFLOPS", gflops_pcie);
        resultDB.AddResult(testName+"_Parity", "cubic", "n", totalTransferTime / total );
    }


    // Release Memory
    err = clReleaseMemObject(d_t);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_p);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_y);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_wdot);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_rf);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_rb);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_c);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_eg);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_rklow);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_a);
    CL_CHECK_ERROR(err);

    // Cleanup Host Memory Objects
    delete[] h_t;
    delete[] h_p;
    delete[] h_y;
    delete[] h_wdot;
    delete[] h_molwt;

}
예제 #25
0
파일: Sort.cpp 프로젝트: Poojachitral/shoc
void runTest(const string& testName, cl_device_id dev, cl_context ctx,
        cl_command_queue queue, ResultDatabase& resultDB, OptionParser& op,
        const string& compileFlags)
{
    int err = 0;

    // Program Setup
    cl_program prog = clCreateProgramWithSource(ctx, 
                                                1, 
                                                &cl_source_sort,
                                                NULL,
                                                &err);
    CL_CHECK_ERROR(err);

    // Before proceeding, make sure the kernel code compiles and
    // all kernels are valid.
    cout << "Compiling sort kernels." << endl;
    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    CL_CHECK_ERROR(err);

    if (err != CL_SUCCESS)
    {
        char log[5000];
        size_t retsize = 0;
        err = clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 5000
                * sizeof(char), log, &retsize);

        CL_CHECK_ERROR(err);
        cout << "Build error." << endl;
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

    // Extract out the 3 kernels
    // Note that these kernels are analogs of those in use for
    // scan, but have had "visiting" logic added to them
    // as described by Merrill et al. See
    // http://www.cs.virginia.edu/~dgm4d/
    cl_kernel reduce = clCreateKernel(prog, "reduce", &err);
    CL_CHECK_ERROR(err);

    cl_kernel top_scan = clCreateKernel(prog, "top_scan", &err);
    CL_CHECK_ERROR(err);
    
    cl_kernel bottom_scan = clCreateKernel(prog, "bottom_scan", &err);
    CL_CHECK_ERROR(err);

    // If the device doesn't support at least 256 work items in a
    // group, use a different kernel (TODO)
    if (getMaxWorkGroupSize(dev) < 256)
    {
        cout << "Scan requires work group size of at least 256" << endl;
        char atts[1024] = "GSize_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (int k = 0; k < passes; k++) {
            resultDB.AddResult(testName , atts, "GB/s", FLT_MAX);
            resultDB.AddResult(testName+"_PCIe" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult(testName+"_Parity" , atts, "GB/s", FLT_MAX);
        }
        return;
    }

    // Problem Sizes
    int probSizes[4] = { 1, 8, 32, 64 };
    int size = probSizes[op.getOptionInt("size")-1];

    // Convert to MiB
    size = (size * 1024 * 1024) / sizeof(T);

    // Create input data on CPU
    unsigned int bytes = size * sizeof(T);

    // Allocate pinned host memory for input data (h_idata)
    cl_mem h_i = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_idata = (T*)clEnqueueMapBuffer(queue, h_i, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);
    
    // Allocate pinned host memory for output data (h_odata)
    cl_mem h_o = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_odata = (T*)clEnqueueMapBuffer(queue, h_o, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);

    // Initialize host memory
    cout << "Initializing host memory." << endl;
    for (int i = 0; i < size; i++)
    {
        h_idata[i] = i % 16; // Fill with some pattern
        h_odata[i] = -1;
    }

    // The radix width in bits
    const int radix_width = 4; // Changing this requires major kernel updates
    const int num_digits = (int)pow((double)2, radix_width); // n possible digits

    // Allocate device memory for input array
    cl_mem d_idata = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    
    // Allocate device memory for output array
    cl_mem d_odata = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
    CL_CHECK_ERROR(err);

    // Number of local work items per group
    const size_t local_wsize  = 256;
    
    // Number of global work items
    const size_t global_wsize = 16384; // i.e. 64 work groups
    const size_t num_work_groups = global_wsize / local_wsize;

    // Allocate device memory for local work group intermediate sums
    cl_mem d_isums = clCreateBuffer(ctx, CL_MEM_READ_WRITE, 
            num_work_groups * num_digits * sizeof(T), NULL, &err);
    CL_CHECK_ERROR(err);

    // Set the kernel arguments for the reduction kernel
    err = clSetKernelArg(reduce, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 1, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 2, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 3, local_wsize * sizeof(T), NULL);
    CL_CHECK_ERROR(err);
                        
    // Set the kernel arguments for the top-level scan
    err = clSetKernelArg(top_scan, 0, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(top_scan, 1, sizeof(cl_int), (void*)&num_work_groups);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(top_scan, 2, local_wsize * 2 * sizeof(T), NULL);
    CL_CHECK_ERROR(err);

    // Set the kernel arguments for the bottom-level scan
    err = clSetKernelArg(bottom_scan, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 1, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 2, sizeof(cl_mem), (void*)&d_odata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 3, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 4, local_wsize * 2 * sizeof(T), NULL);
    CL_CHECK_ERROR(err);
    
    // Copy data to GPU
    cout << "Copying input data to device." << endl;
    Event evTransfer("PCIe transfer");
    err = clEnqueueWriteBuffer(queue, d_idata, true, 0, bytes, h_idata, 0,
            NULL, &evTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    err = clFinish(queue);
    CL_CHECK_ERROR(err);
    evTransfer.FillTimingInfo();
    double inTransferTime = evTransfer.StartEndRuntime();

    // Repeat the test multiplie times to get a good measurement
    int passes = op.getOptionInt("passes");

    cout << "Running benchmark with size " << size << endl;
    for (int k = 0; k < passes; k++)
    {
        int th = Timer::Start();
        // Assuming an 8 bit byte.
        for (int shift = 0; shift < sizeof(T)*8; shift += radix_width)
        {
            // Like scan, we use a reduce-then-scan approach
            
            // But before proceeding, update the shift appropriately
            // for each kernel. This is how many bits to shift to the
            // right used in binning.
            err = clSetKernelArg(reduce, 4, sizeof(cl_int), (void*)&shift);
            CL_CHECK_ERROR(err);

            err = clSetKernelArg(bottom_scan, 5, sizeof(cl_int), (void*)&shift);
            CL_CHECK_ERROR(err);
            
            // Also, the sort is not in place, so swap the input and output
            // buffers on each pass.
            bool even = ((shift / radix_width) % 2 == 0) ? true : false;
            
            if (even)
            {
                // Set the kernel arguments for the reduction kernel
                err = clSetKernelArg(reduce, 0, sizeof(cl_mem), 
                        (void*)&d_idata);
                CL_CHECK_ERROR(err);
                // Set the kernel arguments for the bottom-level scan
                err = clSetKernelArg(bottom_scan, 0, sizeof(cl_mem), 
                        (void*)&d_idata);
                CL_CHECK_ERROR(err);
                err = clSetKernelArg(bottom_scan, 2, sizeof(cl_mem), 
                        (void*)&d_odata);
                CL_CHECK_ERROR(err);
            } 
            else // i.e. odd pass
            {
                // Set the kernel arguments for the reduction kernel
                err = clSetKernelArg(reduce, 0, sizeof(cl_mem), 
                        (void*)&d_odata);
                CL_CHECK_ERROR(err);
                // Set the kernel arguments for the bottom-level scan
                err = clSetKernelArg(bottom_scan, 0, sizeof(cl_mem), 
                        (void*)&d_odata);
                CL_CHECK_ERROR(err);
                err = clSetKernelArg(bottom_scan, 2, sizeof(cl_mem), 
                        (void*)&d_idata);
                CL_CHECK_ERROR(err);
            }

            // Each thread block gets an equal portion of the
            // input array, and computes occurrences of each digit.
            err = clEnqueueNDRangeKernel(queue, reduce, 1, NULL,
                        &global_wsize, &local_wsize, 0, NULL, NULL);
            
            // Next, a top-level exclusive scan is performed on the
            // per block histograms.  This is done by a single
            // work group (note global size here is the same as local).
            err = clEnqueueNDRangeKernel(queue, top_scan, 1, NULL,
                        &local_wsize, &local_wsize, 0, NULL, NULL);
          
            // Finally, a bottom-level scan is performed by each block
            // that is seeded with the scanned histograms which rebins,
            // locally scans, then scatters keys to global memory
            err = clEnqueueNDRangeKernel(queue, bottom_scan, 1, NULL,
                        &global_wsize, &local_wsize, 0, NULL, NULL);
        }
        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        double total_sort = Timer::Stop(th, "total sort time");

        err = clEnqueueReadBuffer(queue, d_idata, true, 0, bytes, h_odata,
                0, NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        double totalTransfer = inTransferTime + evTransfer.StartEndRuntime();
        totalTransfer /= 1.e9; // Convert to seconds

        // If answer is incorrect, stop test and do not report performance
        if (! verifySort(h_odata, size))
        {
            return;
        }
        
        char atts[1024];
        double avgTime = total_sort;
        double gbs = (double) (size * sizeof(T)) / (1000. * 1000. * 1000.);
        sprintf(atts, "%d_items", size);
        resultDB.AddResult(testName, atts, "GB/s", gbs / (avgTime));
        resultDB.AddResult(testName+"_PCIe", atts, "GB/s",
                gbs / (avgTime + totalTransfer));
        resultDB.AddResult(testName+"_Parity", atts, "N",
                totalTransfer / avgTime);
    }

    // Clean up device memory
    err = clReleaseMemObject(d_idata);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_odata);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_isums);
    CL_CHECK_ERROR(err);

    // Clean up pinned host memory
    err = clEnqueueUnmapMemObject(queue, h_i, h_idata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_o, h_odata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_i);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_o);
    CL_CHECK_ERROR(err);
    
    // Clean up program and kernel objects
    err = clReleaseProgram(prog);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(reduce);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(top_scan);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(bottom_scan);
    CL_CHECK_ERROR(err);
}
예제 #26
0
파일: tpScan.cpp 프로젝트: dylanzika/shoc
void runTest(const string& testName, cl_device_id dev, cl_context ctx,
        cl_command_queue queue, ResultDatabase& resultDB, OptionParser& op,
        const string& compileFlags)
{
    int err = 0;

    // Collect basic MPI information
    int mpi_size, mpi_rank;
    MPI_Comm_size(MPI_COMM_WORLD, &mpi_size);
    MPI_Comm_rank(MPI_COMM_WORLD, &mpi_rank);

    // Program Setup
    cl_program prog = clCreateProgramWithSource(ctx,
                                                1,
                                                &cl_source_scan,
                                                NULL,
                                                &err);
    CL_CHECK_ERROR(err);

    // Before proceeding, make sure the kernel code compiles and
    // all kernels are valid.
    if (mpi_rank == 0)
    {
        cout << "Compiling scan kernels." << endl;
    }
    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    CL_CHECK_ERROR(err);

    if (err != CL_SUCCESS)
    {
        char log[5000];
        size_t retsize = 0;
        err = clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 5000
                * sizeof(char), log, &retsize);

        CL_CHECK_ERROR(err);
        cout << "Build error." << endl;
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

    // Extract out the 3 kernels
    cl_kernel reduce = clCreateKernel(prog, "reduce", &err);
    CL_CHECK_ERROR(err);

    cl_kernel top_scan = clCreateKernel(prog, "top_scan", &err);
    CL_CHECK_ERROR(err);

    cl_kernel bottom_scan = clCreateKernel(prog, "bottom_scan", &err);
    CL_CHECK_ERROR(err);

    // If the device doesn't support at least 256 work items in a
    // group, use a different kernel (TODO)
    if (getMaxWorkGroupSize(dev) < 256)
    {
        cout << "Scan requires work group size of at least 256" << endl;
        char atts[1024] = "GSize_Not_Supported";
        // resultDB requires neg entry for every possible result
        int passes = op.getOptionInt("passes");
        for (int k = 0; k < passes; k++)
        {
            resultDB.AddResult(testName+"-Kernel" , atts, "GB/s", FLT_MAX);
            resultDB.AddResult(testName+"-Kernel+PCIe" , atts, "GB/s",
                FLT_MAX);
            resultDB.AddResult(testName+"-MPI_ExScan" , atts, "GB/s",
                FLT_MAX);
            resultDB.AddResult(testName+"-Overall" , atts, "GB/s", FLT_MAX);
        }
        return;
    }

    // Problem Sizes
    int probSizes[4] = { 1, 8, 32, 64 };
    int size = probSizes[op.getOptionInt("size")-1];

    // Convert to MB
    size = (size * 1024 * 1024) / sizeof(T);

    // Create input data on CPU
    unsigned int bytes = size * sizeof(T);

    // Allocate pinned host memory for input data (h_idata)
    cl_mem h_i = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_idata = (T*)clEnqueueMapBuffer(queue, h_i, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);

    // Allocate pinned host memory for output data (h_odata)
    cl_mem h_o = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            bytes, NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_odata = (T*)clEnqueueMapBuffer(queue, h_o, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, bytes, 0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);


    // Initialize host memory
    if (mpi_rank == 0)
    {
        cout << "Initializing host memory." << endl;
    }
    for (int i = 0; i < size; i++)
    {
        h_idata[i] = i % 2; //Fill with some pattern
        h_odata[i] = -1;
    }

    // Allocate device memory for input array
    cl_mem d_idata = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
    CL_CHECK_ERROR(err);

    // Allocate device memory for output array
    cl_mem d_odata = clCreateBuffer(ctx, CL_MEM_READ_WRITE, bytes, NULL, &err);
    CL_CHECK_ERROR(err);

    // Number of local work items per group
    const size_t local_wsize  = 256;

    // Number of local work groups and total work items
    const size_t num_work_groups = 64;
    const size_t global_wsize = local_wsize * num_work_groups;

    // Allocate device memory for local work group intermediate sums
    cl_mem d_isums = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
            num_work_groups * sizeof(T), NULL, &err);
    CL_CHECK_ERROR(err);

    // Allocate pinned host memory for intermediate block sums (h_isums)
    cl_mem h_b = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
        num_work_groups * sizeof(T), NULL, &err);
    CL_CHECK_ERROR(err);
    T* h_isums = (T*)clEnqueueMapBuffer(queue, h_b, true,
        CL_MAP_READ|CL_MAP_WRITE, 0, num_work_groups * sizeof(T),
        0, NULL, NULL, &err);
    CL_CHECK_ERROR(err);

    // Set the kernel arguments for the reduction kernel
    err = clSetKernelArg(reduce, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 1, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 2, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(reduce, 3, local_wsize * sizeof(T), NULL);
    CL_CHECK_ERROR(err);

    // Set the kernel arguments for the top-level scan
    err = clSetKernelArg(top_scan, 0, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(top_scan, 1, sizeof(cl_int), (void*)&num_work_groups);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(top_scan, 2, local_wsize * 2 * sizeof(T), NULL);
    CL_CHECK_ERROR(err);

    // Set the kernel arguments for the bottom-level scan
    err = clSetKernelArg(bottom_scan, 0, sizeof(cl_mem), (void*)&d_idata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 1, sizeof(cl_mem), (void*)&d_isums);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 2, sizeof(cl_mem), (void*)&d_odata);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 3, sizeof(cl_int), (void*)&size);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(bottom_scan, 4, local_wsize * 2 * sizeof(T), NULL);
    CL_CHECK_ERROR(err);

    // Repeat the test multiple times to get a good measurement
    int passes = op.getOptionInt("passes");

    if (mpi_rank == 0)
    {
        cout << "Running benchmark with size " << size << endl;
    }
    for (int k = 0; k < passes; k++)
    {
        // Timing variables
        double pcie_time=0., kernel_time=0., mpi_time=0.;

        // Copy data to GPU
        Event evTransfer("PCIe transfer");
        double time_temp = 0.;
        err = clEnqueueWriteBuffer(queue, d_idata, true, 0, bytes, h_idata, 0,
                NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        pcie_time += (double)evTransfer.StartEndRuntime() / 1e9;

        // This code uses a reduce-then-scan strategy.
        // The major steps of the algorithm are:
        // 1. Local reduction on a node
        // 2. Global exclusive scan of the reduction values
        // 3. Local inclusive scan, seeded with the node's result
        //    from the global exclusive scan
        Event ev_reduce("Reduction Kernel");
        err = clEnqueueNDRangeKernel(queue, reduce, 1, NULL,
                    &global_wsize, &local_wsize, 0, NULL,
                    &ev_reduce.CLEvent());
        err = clFinish(queue);
        ev_reduce.FillTimingInfo();
        kernel_time += (double)ev_reduce.StartEndRuntime() * 1e-9;

        // Next step is to copy the reduced blocks back to the host,
        // sum them, and perform the MPI exlcusive (top level) scan.
        err = clEnqueueReadBuffer(queue, d_isums, true, 0,
                num_work_groups*sizeof(T), h_isums, 0,
                NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        pcie_time += (double)evTransfer.StartEndRuntime() * 1e-9;

        // Start the timer for MPI Scan
        int globscan_th = Timer::Start();
        T reduced=0., scanned=0.;

        // To get the true sum for this node, we have to add up
        // the block sums before MPI scanning.
        for (int i = 0; i < num_work_groups; i++)
        {
            reduced += h_isums[i];
        }

        // Next step is an exclusive scan across MPI ranks.
        // Then a local scan seeded with the result from MPI.
        globalExscan(&reduced, &scanned);
        mpi_time += Timer::Stop(globscan_th, "Global Scan");

        // Now, scanned contains all the information we need from other nodes
        // Next step is to perform the local top level (i.e. across blocks) scan,
        // but seed it with the "scanned", the sum of elems on all lower ranks.
        h_isums[0] += scanned;

        err = clEnqueueWriteBuffer(queue, d_isums, true, 0, sizeof(T), h_isums, 0,
                NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);
        evTransfer.FillTimingInfo();
        pcie_time += (double)evTransfer.StartEndRuntime() * 1e-9;

        Event ev_scan("Scan Kernel");
        err = clEnqueueNDRangeKernel(queue, top_scan, 1, NULL,
                &local_wsize, &local_wsize, 0, NULL, &ev_scan.CLEvent());
        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        ev_scan.FillTimingInfo();
        kernel_time += ((double)ev_scan.StartEndRuntime() * 1.e-9);

        // Finally, a bottom-level scan is performed by each block
        // that is seeded with the scanned value in block sums
        err = clEnqueueNDRangeKernel(queue, bottom_scan, 1, NULL,
                    &global_wsize, &local_wsize, 0, NULL,
                    &ev_scan.CLEvent());
        err = clFinish(queue);
        CL_CHECK_ERROR(err);
        ev_scan.FillTimingInfo();
        kernel_time += ((double)ev_scan.StartEndRuntime() * 1.e-9);

        // Read data back for correctness check
        err = clEnqueueReadBuffer(queue, d_odata, true, 0, bytes, h_odata,
                0, NULL, &evTransfer.CLEvent());
        CL_CHECK_ERROR(err);

        // Lightweight correctness check -- won't apply
        // if data is not initialized to i%2 above
        if (mpi_rank == mpi_size-1)
        {
            if (h_odata[size-1] != (mpi_size * size) / 2)
            {
                cout << "Test Failed\n";
            }
            else
            {
                cout << "Test Passed\n";
            }
        }

        char atts[1024];
        sprintf(atts, "%d items", size);
        double global_gb = (double)(mpi_size * size * sizeof(T)) /
            (1000. * 1000. * 1000.);

        resultDB.AddResult(testName+"-Kernel" , atts, "GB/s",
                global_gb / kernel_time);
        resultDB.AddResult(testName+"-Kernel+PCIe" , atts, "GB/s",
                global_gb / (kernel_time + pcie_time));
        resultDB.AddResult(testName+"-MPI_ExScan" , atts, "GB/s",
                (mpi_size * sizeof(T) *1e-9) / mpi_time);
        resultDB.AddResult(testName+"-Overall" , atts, "GB/s",
                global_gb / (kernel_time + pcie_time + mpi_time));
    }

    // Clean up device memory
    err = clReleaseMemObject(d_idata);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_odata);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_isums);
    CL_CHECK_ERROR(err);

    // Clean up pinned host memory
    err = clEnqueueUnmapMemObject(queue, h_i, h_idata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_o, h_odata, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_b, h_isums, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_i);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_o);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_b);
    CL_CHECK_ERROR(err);

    err = clReleaseProgram(prog);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(reduce);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(top_scan);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(bottom_scan);
    CL_CHECK_ERROR(err);
}
예제 #27
0
파일: Spmv.cpp 프로젝트: bart-utahman/shoc
void ellPackTest(cl_device_id dev, cl_context ctx, string compileFlags,
                 cl_command_queue queue, ResultDatabase& resultDB,
                 OptionParser& op, floatType* h_val, int* h_cols,
                 int* h_rowDelimiters, floatType* h_vec, floatType* h_out,
                 int numRows, int numNonZeroes, floatType* refOut, bool padded,
                 int paddedSize, const size_t maxImgWidth)
{
    if (devSupportsImages)
    {
        char texflags[64];
        sprintf(texflags," -DUSE_TEXTURE -DMAX_IMG_WIDTH=%ld", maxImgWidth);
        compileFlags+=string(texflags);
    }

    // Set up OpenCL Program Object
    int err = 0;
    cl_program prog = clCreateProgramWithSource(ctx, 1, &cl_source_spmv, NULL,
            &err);
    CL_CHECK_ERROR(err);

    // Build the openCL kernels
    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    CL_CHECK_ERROR(err);

    // If there is a build error, print the output and return
    if (err != CL_SUCCESS)
    {
        char log[5000];
        size_t retsize = 0;
        err = clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 50000
                * sizeof(char), log, &retsize);
        CL_CHECK_ERROR(err);
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

    int *h_rowLengths = new int[paddedSize];
    int maxrl = 0;
    for (int k=0; k<numRows; k++)
    {
        h_rowLengths[k] = h_rowDelimiters[k+1] - h_rowDelimiters[k];
        if (h_rowLengths[k] > maxrl)
        {
            maxrl = h_rowLengths[k];
        }
    }
    for (int p=numRows; p < paddedSize; p++)
    {
        h_rowLengths[p] = 0;
    }

    // Column major format host data structures
    int cmSize = padded ? paddedSize : numRows;
    floatType *h_valcm = new floatType[maxrl * cmSize];
    int *h_colscm = new int[maxrl * cmSize];
    convertToColMajor(h_val, h_cols, numRows, h_rowDelimiters, h_valcm,
                              h_colscm, h_rowLengths, maxrl, padded);

    // Device data structures
    cl_mem d_val, d_vec, d_out; // floating point
    cl_mem d_cols, d_rowLengths; // integer

    // Allocate device memory
    d_val = clCreateBuffer(ctx, CL_MEM_READ_WRITE, maxrl * cmSize *
        sizeof(clFloatType), NULL, &err);
    CL_CHECK_ERROR(err);
    d_cols = clCreateBuffer(ctx, CL_MEM_READ_WRITE, maxrl * cmSize *
        sizeof(int), NULL, &err);
    CL_CHECK_ERROR(err);
    int imgHeight = 0;
    if (devSupportsImages)
    {
        imgHeight=(numRows+maxImgWidth-1)/maxImgWidth;
        cl_image_format fmt; fmt.image_channel_data_type=CL_FLOAT;
        if(sizeof(floatType)==4)
        fmt.image_channel_order=CL_R;
        else
        fmt.image_channel_order=CL_RG;
        d_vec = clCreateImage2D( ctx, CL_MEM_READ_ONLY, &fmt, maxImgWidth,
            imgHeight, 0, NULL, &err);
        CL_CHECK_ERROR(err);
    } else {
        d_vec = clCreateBuffer(ctx, CL_MEM_READ_WRITE, numRows *
            sizeof(clFloatType), NULL, &err);
        CL_CHECK_ERROR(err);
    }
    d_out = clCreateBuffer(ctx, CL_MEM_READ_WRITE, paddedSize *
        sizeof(clFloatType), NULL, &err);
    CL_CHECK_ERROR(err);
    d_rowLengths = clCreateBuffer(ctx, CL_MEM_READ_WRITE, cmSize *
        sizeof(int), NULL, &err);
    CL_CHECK_ERROR(err);

    // Setup events for timing
    Event valTransfer("transfer Val data over PCIe bus");
    Event colsTransfer("transfer cols data over PCIe bus");
    Event vecTransfer("transfer vec data over PCIe bus");
    Event rowLengthsTransfer("transfer rowLengths data over PCIe bus");

    // Transfer data to device
    err = clEnqueueWriteBuffer(queue, d_val, true, 0, maxrl * cmSize *
        sizeof(clFloatType), h_valcm, 0, NULL, &valTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    err = clEnqueueWriteBuffer(queue, d_cols, true, 0, maxrl * cmSize *
        sizeof(cl_int), h_colscm, 0, NULL, &colsTransfer.CLEvent());
    CL_CHECK_ERROR(err);

    if (devSupportsImages)
    {
        size_t offset[3]={0};
        size_t size[3]={maxImgWidth,(size_t)imgHeight,1};
        err = clEnqueueWriteImage(queue,d_vec, true, offset, size,
            0, 0, h_vec, 0, NULL, &vecTransfer.CLEvent());
        CL_CHECK_ERROR(err);
    } else {
        err = clEnqueueWriteBuffer(queue, d_vec, true, 0, numRows *
            sizeof(clFloatType), h_vec, 0, NULL, &vecTransfer.CLEvent());
        CL_CHECK_ERROR(err);
    }

    err = clEnqueueWriteBuffer(queue, d_rowLengths, true, 0, cmSize *
        sizeof(int), h_rowLengths, 0, NULL, &rowLengthsTransfer.CLEvent());
    CL_CHECK_ERROR(err);

    err = clFinish(queue);
    CL_CHECK_ERROR(err);

    valTransfer.FillTimingInfo();
    colsTransfer.FillTimingInfo();
    vecTransfer.FillTimingInfo();
    rowLengthsTransfer.FillTimingInfo();

    double iTransferTime =  valTransfer.StartEndRuntime() +
                           colsTransfer.StartEndRuntime() +
                            vecTransfer.StartEndRuntime() +
                     rowLengthsTransfer.StartEndRuntime();

    // Set up kernel arguments
    cl_kernel ellpackr = clCreateKernel(prog, "spmv_ellpackr_kernel", &err);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 0, sizeof(cl_mem), (void*) &d_val);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 1, sizeof(cl_mem), (void*) &d_vec);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 2, sizeof(cl_mem), (void*) &d_cols);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 3, sizeof(cl_mem), (void*) &d_rowLengths);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 4, sizeof(cl_int), (void*) &cmSize);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(ellpackr, 5, sizeof(cl_mem), (void*) &d_out);
    CL_CHECK_ERROR(err);

    const size_t globalWorkSize = cmSize;
    const size_t localWorkSize = BLOCK_SIZE;
    Event kernelExec("ELLPACKR Kernel Execution");

    int passes = op.getOptionInt("passes");
    int iters  = op.getOptionInt("iterations");

    for (int k = 0; k < passes; k++)
    {
        double totalKernelTime = 0.0;
        for (int j = 0; j < iters; j++)
        {
            err = clEnqueueNDRangeKernel(queue, ellpackr, 1, NULL,
                &globalWorkSize, &localWorkSize, 0, NULL,
                &kernelExec.CLEvent());
            CL_CHECK_ERROR(err);
            err = clFinish(queue);
            CL_CHECK_ERROR(err);
            kernelExec.FillTimingInfo();
            totalKernelTime += kernelExec.StartEndRuntime();
        }

         Event outTransfer("d->h data transfer");
         err = clEnqueueReadBuffer(queue, d_out, true, 0, numRows *
             sizeof(clFloatType), h_out, 0, NULL, &outTransfer.CLEvent());
         CL_CHECK_ERROR(err);
         err = clFinish(queue);
         CL_CHECK_ERROR(err);
         outTransfer.FillTimingInfo();
         double oTransferTime = outTransfer.StartEndRuntime();

        // Compare reference solution to GPU result
        if (! verifyResults(refOut, h_out, numRows, k)) {
            return;
        }
        char atts[TEMP_BUFFER_SIZE];
        char benchName[TEMP_BUFFER_SIZE];
        double avgTime = totalKernelTime / (double)iters;
        sprintf(atts, "%d_elements_%d_rows", numNonZeroes, cmSize);
        double gflop = 2 * (double) numNonZeroes;
        bool dpTest = (sizeof(floatType) == sizeof(double));
        sprintf(benchName, "%sELLPACKR-%s", padded ? "Padded_":"",
                dpTest ? "DP":"SP");
        resultDB.AddResult(benchName, atts, "Gflop/s", gflop/avgTime);
        sprintf(benchName, "%s_PCIe", benchName);
        resultDB.AddResult(benchName, atts, "Gflop/s", gflop /
            (avgTime + iTransferTime + oTransferTime));
    }

    err = clReleaseProgram(prog);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(ellpackr);
    CL_CHECK_ERROR(err);

    // Free device memory
    err = clReleaseMemObject(d_rowLengths);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_vec);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_out);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_val);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_cols);
    CL_CHECK_ERROR(err);

    // Free host memory
    delete[] h_rowLengths;
    delete[] h_valcm;
    delete[] h_colscm;
}
예제 #28
0
void runTest(const string& testName, cl_device_id dev, cl_context ctx,
        cl_command_queue queue, ResultDatabase& resultDB, OptionParser& op,
        string compileFlags)
{
    // Problem Parameters
    const int probSizes[4] = { 12288, 24576, 36864, 73728 };
    int sizeClass = op.getOptionInt("size");
    assert(sizeClass >= 0 && sizeClass < 5);
    int nAtom = probSizes[sizeClass - 1];

    // Allocate problem data on host
    cl_mem h_pos, h_force, h_neigh;
    posVecType*   position;
    forceVecType* force;
    int* neighborList;
    int passes = op.getOptionInt("passes");
    int iter   = op.getOptionInt("iterations");

    // Allocate and map pinned host memory
    int err = 0;
    // Position
    h_pos = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            sizeof(posVecType)*nAtom, NULL, &err);
    CL_CHECK_ERROR(err);
    position = (posVecType*)clEnqueueMapBuffer(queue, h_pos, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, sizeof(posVecType)*nAtom , 0,
            NULL, NULL, &err);
    CL_CHECK_ERROR(err);
    // Force
    h_force = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            sizeof(forceVecType)*nAtom, NULL, &err);
    CL_CHECK_ERROR(err);
    force = (forceVecType*)clEnqueueMapBuffer(queue, h_force, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, sizeof(forceVecType)*nAtom , 0,
            NULL, NULL, &err);
    CL_CHECK_ERROR(err);
    // Neighbor List
    h_neigh = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
            sizeof(int) * nAtom * maxNeighbors, NULL, &err);
    CL_CHECK_ERROR(err);
    neighborList = (int*)clEnqueueMapBuffer(queue, h_neigh, true,
            CL_MAP_READ|CL_MAP_WRITE, 0, sizeof(int) * nAtom * maxNeighbors, 0,
            NULL, NULL, &err);
    CL_CHECK_ERROR(err);

    // Allocate device memory
    cl_mem d_force = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
            nAtom * sizeof(forceVecType), NULL, &err);
    CL_CHECK_ERROR(err);
    cl_mem d_position = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
            nAtom * sizeof(posVecType), NULL, &err);
    CL_CHECK_ERROR(err);
    // Allocate device memory neighbor list
    cl_mem d_neighborList = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
            maxNeighbors * nAtom * sizeof(int), NULL, &err);
    CL_CHECK_ERROR(err);

    size_t maxGroupSize = getMaxWorkGroupSize(dev);
    if (maxGroupSize < 128)
    {
        cout << "MD requires a work group size of at least 128" << endl;
        // Add special values to the results database
        char atts[1024];
        sprintf(atts, "GSize_Not_Supported");
        for (int i=0 ; i<passes ; ++i) {
            resultDB.AddResult(testName, atts, "GFLOPS", FLT_MAX);
            resultDB.AddResult(testName + "_PCIe", atts, "GFLOPS", FLT_MAX);
            resultDB.AddResult(testName+"-Bandwidth", atts, "GB/s", FLT_MAX);
            resultDB.AddResult(testName+"-Bandwidth_PCIe", atts, "GB/s", FLT_MAX);
            resultDB.AddResult(testName+"_Parity", atts, "N", FLT_MAX);
        }
        return;
    }
    size_t localSize  = 128;
    size_t globalSize = nAtom;

    cout << "Initializing test problem (this can take several "
            "minutes for large problems).\n                   ";

    // Seed random number generator
    srand48(8650341L);

    // Initialize positions -- random distribution in cubic domain
    for (int i = 0; i < nAtom; i++)
    {
        position[i].x = (drand48() * domainEdge);
        position[i].y = (drand48() * domainEdge);
        position[i].z = (drand48() * domainEdge);
    }

    // Copy position to GPU
    Event evTransfer("h->d transfer");
    err = clEnqueueWriteBuffer(queue, d_position, true, 0,
            nAtom * sizeof(posVecType), position, 0, NULL,
            &evTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    err = clFinish(queue);
    CL_CHECK_ERROR(err);

    evTransfer.FillTimingInfo();
    long transferTime = evTransfer.StartEndRuntime();

    // Keep track of how many atoms are within the cutoff distance to
    // accurately calculate FLOPS later
    int totalPairs = buildNeighborList<T, posVecType>(nAtom, position,
            neighborList);

    cout << "Finished.\n";
    cout << totalPairs << " of " << nAtom*maxNeighbors <<
            " pairs within cutoff distance = " <<
            100.0 * ((double)totalPairs / (nAtom*maxNeighbors)) << " %" << endl;

    // Copy data to GPU
    err = clEnqueueWriteBuffer(queue, d_neighborList, true, 0,
            maxNeighbors * nAtom * sizeof(int), neighborList, 0, NULL,
            &evTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    clFinish(queue);
    evTransfer.FillTimingInfo();
    transferTime += evTransfer.StartEndRuntime();

    // Build the openCL kernel
    cl_program prog = clCreateProgramWithSource(ctx, 1, &cl_source_md, NULL,
            &err);
    CL_CHECK_ERROR(err);

    err = clBuildProgram(prog, 1, &dev, compileFlags.c_str(), NULL, NULL);
    CL_CHECK_ERROR(err);

    // If there is a build error, print the output and return
    if (err != CL_SUCCESS)
    {
        char log[5000];
        size_t retsize = 0;
        err = clGetProgramBuildInfo(prog, dev, CL_PROGRAM_BUILD_LOG, 50000
                * sizeof(char), log, &retsize);
        CL_CHECK_ERROR(err);
        cout << "Retsize: " << retsize << endl;
        cout << "Log: " << log << endl;
        return;
    }

    // Extract out the kernels
    cl_kernel lj_kernel = clCreateKernel(prog, "compute_lj_force",
            &err);
    CL_CHECK_ERROR(err);

    T lj1_t = (T) lj1;
    T lj2_t = (T) lj2;
    T cutsq_t = (T) cutsq;

    // Set kernel arguments
    err = clSetKernelArg(lj_kernel, 0, sizeof(cl_mem),
            (void*) &d_force);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 1, sizeof(cl_mem),
            (void*) &d_position);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 2, sizeof(cl_int),
            (void*) &maxNeighbors);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 3, sizeof(cl_mem),
            (void*) &d_neighborList);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 4, sizeof(T),
            (void*) &cutsq_t);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 5, sizeof(T),
            (void*) &lj1_t);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 6, sizeof(T),
            (void*) &lj2_t);
    CL_CHECK_ERROR(err);
    err = clSetKernelArg(lj_kernel, 7, sizeof(cl_int),
            (void*) &nAtom);
    CL_CHECK_ERROR(err);

    Event evLJ("computeLJ");

    // Warm up the kernel and check correctness
    err = clEnqueueNDRangeKernel(queue, lj_kernel, 1, NULL, &globalSize,
            &localSize, 0, NULL, &evLJ.CLEvent());
    CL_CHECK_ERROR(err);
    err = clFinish(queue);
    CL_CHECK_ERROR(err);

    err = clEnqueueReadBuffer(queue, d_force, true, 0,
            nAtom * sizeof(forceVecType), force, 0, NULL,
            &evTransfer.CLEvent());
    CL_CHECK_ERROR(err);
    err = clFinish(queue);
    CL_CHECK_ERROR(err);
    evTransfer.FillTimingInfo();
    transferTime += evTransfer.StartEndRuntime();


    cout << "Performing Correctness Check (can take several minutes)\n";

    // If results are correct, skip the performance tests
    if (!checkResults<T, forceVecType, posVecType>(force, position,
            neighborList, nAtom))
    {
        return;
    }


    for (int i = 0; i < passes; i++)
    {
        double total_time = 0.0;
        for (int j = 0; j < iter; j++)
        {
            //Launch Kernels
            err = clEnqueueNDRangeKernel(queue, lj_kernel, 1, NULL,
                    &globalSize, &localSize, 0, NULL,
                    &evLJ.CLEvent());
            CL_CHECK_ERROR(err);
            err = clFinish(queue);
            CL_CHECK_ERROR(err);

            // Collect timing info from events
            evLJ.FillTimingInfo();
            total_time += evLJ.SubmitEndRuntime();
        }

        char atts[1024];
        long int nflops = (8 * nAtom * maxNeighbors) + (totalPairs * 13);
        sprintf(atts, "%d_atoms", nAtom);
        total_time /= (double) iter;
        resultDB.AddResult(testName, atts, "GFLOPS",
                ((double) nflops) / total_time);
        resultDB.AddResult(testName + "_PCIe", atts, "GFLOPS",
                        ((double) nflops) / (total_time + transferTime));

        long int numPairs = nAtom * maxNeighbors;
        long int nbytes = (3 * sizeof(T) * (1+numPairs)) + // position data
                          (3 * sizeof(T) * nAtom) + // force for each atom
                          (sizeof(int) * numPairs); // neighbor list
        double gbytes = (double)nbytes / (1000. * 1000. * 1000.);
        double seconds = total_time / 1.e9;
        resultDB.AddResult(testName+"-Bandwidth", atts, "GB/s", gbytes /
                        seconds);
        resultDB.AddResult(testName+"-Bandwidth_PCIe", atts, "GB/s", gbytes /
                           (seconds + (transferTime / 1.e9)));
        resultDB.AddResult(testName+"_Parity", atts, "N",
                                    (transferTime / 1.e9) / seconds);
    }

    // Clean up
    // Host memory
    err = clEnqueueUnmapMemObject(queue, h_pos,   position, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_force, force, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clEnqueueUnmapMemObject(queue, h_neigh, neighborList, 0, NULL, NULL);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_pos);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_force);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(h_neigh);
    CL_CHECK_ERROR(err);

    // Program Objects
    err = clReleaseProgram(prog);
    CL_CHECK_ERROR(err);
    err = clReleaseKernel(lj_kernel);
    CL_CHECK_ERROR(err);

    // Device Memory
    err = clReleaseMemObject(d_force);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_position);
    CL_CHECK_ERROR(err);
    err = clReleaseMemObject(d_neighborList);
    CL_CHECK_ERROR(err);

}
예제 #29
0
파일: Spmv.cpp 프로젝트: bart-utahman/shoc
// ****************************************************************************
// Function: RunBenchmark
//
// Purpose:
//   Executes the sparse matrix - vector multiplication benchmark
//
// Arguments:
//   dev: the opencl device id to use for the benchmark
//   ctx: the opencl context to use for the benchmark
//   queue: the opencl command queue to issue commands to
//   resultDB: stores results from the benchmark
//   op: the options parser / parameter database
//
// Returns:  nothing
// Programmer: Lukasz Wesolowski
// Creation: August 13, 2010
//
// Modifications:
//
// ****************************************************************************
void
RunBenchmark(cl_device_id dev,
                  cl_context ctx,
                  cl_command_queue queue,
                  ResultDatabase &resultDB,
                  OptionParser &op)
{
    //create list of problem sizes
    int probSizes[4] = {1024, 8192, 12288, 16384};
    int sizeClass = op.getOptionInt("size") - 1;

    // Always run single precision test
    // OpenCL doesn't support templated kernels, so we have to use macros
    cout <<"Single precision tests:\n";
    string spMacros = "-DSINGLE_PRECISION ";
    RunTest<float, cl_float>
        (dev, ctx, queue, resultDB, op, spMacros, probSizes[sizeClass]);

    // If double precision is supported, run the DP test
    if (checkExtension(dev, "cl_khr_fp64"))
    {
        cout << "Double precision tests\n";
        string dpMacros = "-DK_DOUBLE_PRECISION ";
        RunTest<double, cl_double>
            (dev, ctx, queue, resultDB, op, dpMacros, probSizes[sizeClass]);
    }
    else if (checkExtension(dev, "cl_amd_fp64"))
    {
        cout << "Double precision tests\n";
        string dpMacros = "-DAMD_DOUBLE_PRECISION ";
        RunTest<double, cl_double>
            (dev, ctx, queue, resultDB, op, dpMacros, probSizes[sizeClass]);
    }
    else
    {
        std::cout << "Double precision not supported by chosen device, skipping" << std::endl;
        // driver script still needs entries for all tests, even if not run
        int nPasses = (int)op.getOptionInt( "passes" );
        for( unsigned int p = 0; p < nPasses; p++ )
        {
            resultDB.AddResult( (const char*)"CSR-Scalar-DP",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"CSR-Scalar-DP_PCIe",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"CSR-Vector-DP",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"CSR-Vector-DP_PCIe",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"ELLPACKR-DP",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"ELLPACKR-DP_PCIe",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"Padded_CSR-Scalar-DP",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"Padded_CSR-Scalar-DP_PCIe",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"Padded_CSR-Vector-DP",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
            resultDB.AddResult( (const char*)"Padded_CSR-Vector-DP_PCIe",
                                "N/A",
                                "Gflop/s",
                                FLT_MAX );
        }
    }
}
예제 #30
0
void runTest(const string& name, ResultDatabase &resultDB, OptionParser& op)
{	
    int i, j;
    void* work, * chk;
    T2* source, * result;
    unsigned long bytes = 0;

    if (op.getOptionInt("MB") == 0) {
        int probSizes[4] = { 1, 8, 96, 256 };
	int sizeIndex = op.getOptionInt("size")-1;
	if (sizeIndex < 0 || sizeIndex >= 4) {
	    cerr << "Invalid size index specified\n";
	    exit(-1);
	}
        bytes = probSizes[sizeIndex];
    } 
    else {
        bytes = op.getOptionInt("MB");
    }

    // Convert to MB
    bytes *= 1024 * 1024;

    bool do_dp = dp<T2>();
    init(op, do_dp);

    int passes = op.getOptionInt("passes");

    // now determine how much available memory will be used
    int half_n_ffts = bytes / (512*sizeof(T2)*2);
    int n_ffts = half_n_ffts * 2;
    int half_n_cmplx = half_n_ffts * 512;
    unsigned long used_bytes = half_n_cmplx * 2 * sizeof(T2);
    double N = half_n_cmplx*2;
    
    // allocate host and device memory
    allocHostBuffer((void**)&source, used_bytes);
    allocHostBuffer((void**)&result, used_bytes);

    // init host memory...
    for (i = 0; i < half_n_cmplx; i++) {
        source[i].x = (rand()/(float)RAND_MAX)*2-1;
        source[i].y = (rand()/(float)RAND_MAX)*2-1;
        source[i+half_n_cmplx].x = source[i].x;
        source[i+half_n_cmplx].y = source[i].y;
    }

    // alloc device memory
    allocDeviceBuffer(&work, used_bytes);
    allocDeviceBuffer(&chk, 1);

    // Copy to device, and record transfer time
    fprintf(stderr, "used_bytes=%d, N=%g\n", used_bytes, N);

    int pcie_TH = Timer::Start();
    copyToDevice(work, source, used_bytes);
    double transfer_time = Timer::Stop(pcie_TH, "PCIe Transfer Time");

    char chk_init = 0;
    copyToDevice(chk, &chk_init, 1);

    const char *sizeStr;
    stringstream ss;
    ss << "N=" << (long)N;
    sizeStr = strdup(ss.str().c_str());

    for (int k=0; k<passes; k++) {
        // time fft kernel
        int TH = Timer::Start();
        forward(work, n_ffts);
        double t = Timer::Stop(TH, "fft");
        double fftsz = 512;
        double Gflops = n_ffts*(5*fftsz*log2(fftsz))/(t*1e9f);
        double gflopsPCIe = n_ffts*(5*fftsz*log2(fftsz)) /
                ((transfer_time+t)*1e9f);
        resultDB.AddResult(name, sizeStr, "GFLOPS", Gflops);
        resultDB.AddResult(name+"_PCIe", sizeStr, "GFLOPS", gflopsPCIe);
        resultDB.AddResult(name+"_Parity", sizeStr, "N", transfer_time / t);

        // time ifft kernel
        TH = Timer::Start();
        inverse(work, n_ffts);
        t = Timer::Stop(TH, "ifft");
        Gflops = n_ffts*(5*fftsz*log2(fftsz))/(t*1e9f);
        gflopsPCIe = n_ffts*(5*fftsz*log2(fftsz)) /
                ((transfer_time+t)*1e9f);
        resultDB.AddResult(name+"-INV", sizeStr, "GFLOPS", Gflops);
        resultDB.AddResult(name+"-INV_PCIe", sizeStr, "GFLOPS", gflopsPCIe);
        resultDB.AddResult(name+"-INV_Parity", sizeStr, "N", transfer_time / t);

        // time check kernel
        int failed = check(work, chk, half_n_ffts, half_n_cmplx);
        cout << "pass " << k << ((failed) ? ": failed\n" : ": passed\n");
    }
    
    freeDeviceBuffer(work);
    freeDeviceBuffer(chk);
    freeHostBuffer(source);
    freeHostBuffer(result);
}