Home My Page Projects Code Snippets Project Openings diderot
Summary Activity Tracker Tasks SCM

SCM Repository

[diderot] View of /branches/vis12-cl/src/compiler/cl-target/fragments/c-run-nobsp.in
ViewVC logotype

View of /branches/vis12-cl/src/compiler/cl-target/fragments/c-run-nobsp.in

Parent Directory Parent Directory | Revision Log Revision Log


Revision 3137 - (download) (annotate)
Thu Mar 26 14:55:52 2015 UTC (4 years, 3 months ago) by jhr
File size: 2568 byte(s)
working on OpenCL scheduling
//! Run the Diderot program (OpenCL version)
//! \param wrld the world-state of the Diderot program
//! \param maxNSteps the limit on the number of super steps; 0 means unlimited
//! \return the number of steps taken.
uint32_t @PREFIX@Run (@PREFIX@World_t *wrld, uint32_t maxNSteps)
{
    cl_int sts;

    if (maxNSteps == 0) maxNSteps = 0xffffffff;  // essentially unlimited

  // iterate until all strands are stable
    uint32_t nSteps = 0, nUpdates = 0;

    double t0 = airTime();

  // work sizes for the program
    size_t globalWorkSize[2];
    size_t localWorkSize[2]; 
    globalWorkSize[0] = wrld->device->numCUs * wrld->device->cuWidth; 
    globalWorkSize[1] = wrld->nWorkers; 
    localWorkSize[0] = wrld->device->cuWidth;
    localWorkSize[1] = 1;

    if (wrld->verboseFlg) {
	fprintf (stderr, "enqueuing UpdateKern for %d x %d x %d workers\n",
	    wrld->device->numCUs, wrld->device->cuWidth, wrld->nWorkers);
    }

    if (((sts = clSetKernelArg (wrld->UpdateKern, 0, sizeof(cl_mem), &(wrld->schedBuf))) != CL_SUCCESS)
    ||  ((sts = clSetKernelArg (wrld->UpdateKern, 1, sizeof(cl_mem), &(wrld->globalsBuf))) != CL_SUCCESS)
    ||  ((sts = clSetKernelArg (wrld->UpdateKern, 2, sizeof(cl_int), &nSteps)) != CL_SUCCESS)
    ||  ((sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->UpdateKern, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL)) != CL_SUCCESS)) {
	ReportOCLError (wrld, sts, "error enqueuing UpdateKern kernel");
	return 0;
    }

    if ((sts = clFinish (wrld->cmdQ)) != CL_SUCCESS) {
	ReportOCLError (wrld, sts, "error finishing UpdateKern kernel");
	return 0;
    }

    t0 = airTime() - t0;
    if (wrld->verboseFlg) {
        fprintf (stderr, "done in %f seconds\n", t0);
    }
    wrld->runTime += t0;

  // get scheduler info
    struct {
      /* the known prefix of the GPU scheduler structure */
	cl_int		nSteps;		// number of executed steps
	cl_int		nStrands;	// total number of allocated strands
	cl_int		nActive;	// number of active strands
	cl_int		nStable;	// number of stable strands
#ifdef DIDEROT_HAS_DIE
	cl_int		nDied;		// number of strands that have died
#endif
    } schedInfo;

    if ((sts = clEnqueueReadBuffer(wrld->cmdQ, wrld->schedBuf, CL_TRUE, 0, sizeof(schedInfo), &schedInfo, 0, 0, 0)) != CL_SUCCESS) {
	ReportOCLError (wrld, sts, "error getting scheduler info");
	exit(1);
    }

    if (wrld->verboseFlg) {
	fprintf(stderr, "Run finished with %d active + %d stable = %d strands\n",
	    schedInfo.nActive, schedInfo.nStable, schedInfo.nStrands);
    }

    return schedInfo.nSteps;

} // @PREFIX@Run

root@smlnj-gforge.cs.uchicago.edu
ViewVC Help
Powered by ViewVC 1.0.0