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

SCM Repository

[diderot] Diff of /branches/pure-cfg/src/lib/cl-target/main.c
ViewVC logotype

Diff of /branches/pure-cfg/src/lib/cl-target/main.c

Parent Directory Parent Directory | Revision Log Revision Log | View Patch Patch

revision 1491, Fri Sep 9 12:28:00 2011 UTC revision 1492, Fri Sep 9 13:13:47 2011 UTC
# Line 16  Line 16 
16  #include <sys/sysctl.h>  #include <sys/sysctl.h>
17  #include <sys/stat.h>  #include <sys/stat.h>
18    
19    // FIXME: the kernels should really be in the lib subdirectory
20    // path to kernels
21    #define KERNELS_FILE_PATH       DIDEROT_INCLUDE_PATH "/cl_kernels/kernels.cl"
22    
23  typedef struct {  typedef struct {
24      cl_int     blkIdx;         // the id of this block      cl_int              blkIdx;         //!< the id of this block
25      cl_int    nActive;        // number of active (status != DIE or STABLE) strands      cl_int              nActive;        //!< number of active (status != DIE or STABLE) strands
26      cl_int    nDead;          // number of strands in the DIE state      cl_int              nDead;          //!< number of strands in the DIE state
27      cl_int    nStabilizing;   // number of new strands in the STABILIZE state      cl_int              nStabilizing;   //!< number of new strands in the STABILIZE state
28      cl_int    nDying;         // number of new strands in the DIE state      cl_int              nDying;         //!< number of new strands in the DIE state
29  } StrandBlock_t;  } StrandBlock_t;
30    
31  typedef struct {  typedef struct {
32      cl_int numStrands;                 // number of strands      cl_int              numStrands;     //!< number of strands
33      cl_int sId;                         // the index into the todo list or queue      cl_int              sId;            //!< the index into the todo list or queue
34      cl_int nextStrand[1];              // index of the next strand to retrieve from the pool      cl_int              nextStrand[1];  //!< index of the next strand to retrieve from the pool
35      cl_int queueSize;                  // number of blocks on the scheduler's queue      cl_int              queueSize;      //!< number of blocks on the scheduler's queue
36      cl_int todoSize;                   // number of blocks on the scheduler's todo list      cl_int              todoSize;       //!< number of blocks on the scheduler's todo list
37      cl_int numAvailable[1];            // number of active strands left to process      cl_int              numAvailable[1]; //!< number of active strands left to process
38  } SchedState_t;  } SchedState_t;
39    
40    
# Line 39  Line 42 
42  static bool     VerboseFlg = false;  static bool     VerboseFlg = false;
43  static bool     TimingFlg = false;  static bool     TimingFlg = false;
44    
 static char * kernelsFile = "cl_kernels/kernels.cl";  
45  #define WARP_SIZE 32  #define WARP_SIZE 32
46    
47  typedef struct {  typedef struct {
# Line 109  Line 111 
111          fprintf (stderr, "initializing globals ...\n");          fprintf (stderr, "initializing globals ...\n");
112      Diderot_InitGlobals ();      Diderot_InitGlobals ();
113    
   /***** FIXME: OpenCL specific stuff goes here.  Things to do:  
    **  
    ** 1) copy data to GPU  
    ** 2) initialize strands  
    ** 3) run strands to termination  
    ** 4) load results from GPU  
    **/  
   
114      Diderot_World_t *wrld = Diderot_Initially ();  // this may not be right for OpenCL      Diderot_World_t *wrld = Diderot_Initially ();  // this may not be right for OpenCL
115    
116      if (! InitCL(clInfo, wrld))      if (! InitCL(clInfo, wrld))
# Line 152  Line 146 
146      Kernel_Args_t kernelArgs;      Kernel_Args_t kernelArgs;
147      SchedState_t scheduler;      SchedState_t scheduler;
148    
   
149      scheduler.numStrands = wrld->numStrands;      scheduler.numStrands = wrld->numStrands;
150      scheduler.nextStrand[0] = 0;      scheduler.nextStrand[0] = 0;
151      scheduler.todoSize = 0;      scheduler.todoSize = 0;
152      scheduler.sId = 0;      scheduler.sId = 0;
153      scheduler.numAvailable[0] = wrld->numStrands;      scheduler.numAvailable[0] = wrld->numStrands;
154    
   
155      size_t globalWorkSize[1] = {0};      size_t globalWorkSize[1] = {0};
156      size_t localWorkSize[1] = {0};      size_t localWorkSize[1] = {0};
157     int dIdx = clInfo->mainDeviceIdx;     int dIdx = clInfo->mainDeviceIdx;
# Line 185  Line 177 
177          strandBlks[i].nDead = 0;          strandBlks[i].nDead = 0;
178          strandBlks[i].nStabilizing = 0;          strandBlks[i].nStabilizing = 0;
179          strandBlks[i].nDying = 0;          strandBlks[i].nDying = 0;
   
180      }      }
181    
182      kernelArgs.schedMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,      kernelArgs.schedMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
# Line 220  Line 211 
211          sizeof(int) * wrld->numStrands, wrld->status, &sts);          sizeof(int) * wrld->numStrands, wrld->status, &sts);
212      CheckErrorCode (sts, "error creating OpenCL world status buffer\n");      CheckErrorCode (sts, "error creating OpenCL world status buffer\n");
213    
   
214      //Setup the Update Kernel's arguments      //Setup the Update Kernel's arguments
215      SetPhase1Args(wrld->kernel[0].kern, &argCount, &kernelArgs);      SetPhase1Args(wrld->kernel[0].kern, &argCount, &kernelArgs);
216      Diderot_LoadGlobals(wrld->context, wrld->kernel[0].kern, wrld->cmdQ, argCount);      Diderot_LoadGlobals(wrld->context, wrld->kernel[0].kern, wrld->cmdQ, argCount);
# Line 242  Line 232 
232    
233          CheckErrorCode (sts, "error in executing update kernel\n");          CheckErrorCode (sts, "error in executing update kernel\n");
234    
235          //FIXME: Reads back the scheduler meta information because the accumlator variable and the queue size        // FIXME: Reads back the scheduler meta information because the accumulator variable and the queue size
236          // needs to be reset for the next interation of the update kernel. As far as I know, you can't do this          // needs to be reset for the next interation of the update kernel. As far as I know, you can't do this
237          // on the GPU side because we don't know when all workgroups are done.          // on the GPU side because we don't know when all workgroups are done.
238          sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),          sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),
239            &scheduler, 0, NULL, NULL);            &scheduler, 0, NULL, NULL);
240            CheckErrorCode (sts, "error in reading back schedule data before compaction kernel\n");
241    
242          scheduler.sId = 0;          scheduler.sId = 0;
243          scheduler.queueSize = 0;          scheduler.queueSize = 0;
   
   
         CheckErrorCode (sts, "error in reading back schedule data before compaction kernel\n");  
   
244          sts = clEnqueueWriteBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),          sts = clEnqueueWriteBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),
245            &scheduler, 0, NULL, NULL);            &scheduler, 0, NULL, NULL);
246          CheckErrorCode (sts, "error in reading back schedule data before compaction kernel\n");          CheckErrorCode (sts, "error in reading back schedule data before compaction kernel\n");
247    
   
248          //Run the compaction kernel on all strands          //Run the compaction kernel on all strands
249          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1].kern, 1, NULL,          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1].kern, 1, NULL,
250              globalWorkSize, localWorkSize, 0, NULL, NULL);              globalWorkSize, localWorkSize, 0, NULL, NULL);
# Line 327  Line 313 
313    
314  static void SetPhase1Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args)  static void SetPhase1Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args)
315  {  {
316        int count = *argCount;
317      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
318    
319      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->inMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->inMem);
320      CheckErrorCode (sts, "error Setting OpenCL strand in-state argument\n");      CheckErrorCode (sts, "error Setting OpenCL strand in-state argument\n");
     *argCount+=1;  
321    
322      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->outMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->outMem);
323      CheckErrorCode (sts, "error Setting OpenCL strand out-state argument\n");      CheckErrorCode (sts, "error Setting OpenCL strand out-state argument\n");
     *argCount+=1;  
   
324    
325      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->statusMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->statusMem);
326      CheckErrorCode (sts, "error Setting OpenCL world status argument\n");      CheckErrorCode (sts, "error Setting OpenCL world status argument\n");
     *argCount+=1;  
327    
328      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->schedMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->schedMem);
329      CheckErrorCode (sts, "error Setting OpenCL scheduler argument\n");      CheckErrorCode (sts, "error Setting OpenCL scheduler argument\n");
     *argCount+=1;  
330    
331      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->blocksMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->blocksMem);
332      CheckErrorCode (sts, "error Setting OpenCL strand blocks argument\n");      CheckErrorCode (sts, "error Setting OpenCL strand blocks argument\n");
     *argCount+=1;  
333    
334      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->strandBlocksIdxsMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->strandBlocksIdxsMem);
335      CheckErrorCode (sts, "error Setting OpenCL strand blocks' indices argument\n");      CheckErrorCode (sts, "error Setting OpenCL strand blocks' indices argument\n");
     *argCount+=1;  
336    
337      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->queueMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->queueMem);
338      CheckErrorCode (sts, "error Setting OpenCL scheduler queue argument\n");      CheckErrorCode (sts, "error Setting OpenCL scheduler queue argument\n");
     *argCount+=1;  
339    
340      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->todoMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->todoMem);
341      CheckErrorCode (sts, "error Setting OpenCL scheduler todo argument\n");      CheckErrorCode (sts, "error Setting OpenCL scheduler todo argument\n");
     *argCount+=1;  
342    
343      sts = clSetKernelArg (kernel,  *(argCount), sizeof(StrandBlock_t), NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(StrandBlock_t), NULL);
344      CheckErrorCode (sts, "error Setting OpenCL local strand block argument\n");      CheckErrorCode (sts, "error Setting OpenCL local strand block argument\n");
345      *argCount+=1;  
346        *argCount = count;
347    
348  }  }
349    
350  static void SetPhase2Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args, int blk_sz)  static void SetPhase2Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args, int blk_sz)
351  {  {
352        int count = *argCount;
353      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
354    
355      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->inMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->inMem);
356      CheckErrorCode (sts, "P2 error Setting OpenCL strand in-state argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL strand in-state argument\n");
     *argCount+=1;  
357    
358      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->outMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->outMem);
359      CheckErrorCode (sts, "P2 error Setting OpenCL strand out-state argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL strand out-state argument\n");
     *argCount+=1;  
360    
361      sts = clSetKernelArg (kernel, *(argCount), sizeof(cl_mem), &args->statusMem);      sts = clSetKernelArg (kernel, count++, sizeof(cl_mem), &args->statusMem);
362      CheckErrorCode (sts, "P2 error Setting OpenCL world status argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL world status argument\n");
     *argCount+=1;  
363    
364      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->schedMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->schedMem);
365      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler argument\n");
     *argCount+=1;  
366    
367      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->blocksMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->blocksMem);
368      CheckErrorCode (sts, "P2 error Setting OpenCL strand blocks argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL strand blocks argument\n");
     *argCount+=1;  
369    
370      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->strandBlocksIdxsMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->strandBlocksIdxsMem);
371      CheckErrorCode (sts, "P2 error Setting OpenCL strand blocks' indices argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL strand blocks' indices argument\n");
     *argCount+=1;  
372    
373      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->queueMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->queueMem);
374      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler queue argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler queue argument\n");
     *argCount+=1;  
375    
376      sts = clSetKernelArg (kernel,  *(argCount), sizeof(cl_mem), &args->todoMem);      sts = clSetKernelArg (kernel,  count++, sizeof(cl_mem), &args->todoMem);
377      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler todo argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL scheduler todo argument\n");
     *argCount+=1;  
378    
379      sts = clSetKernelArg (kernel,  *(argCount), sizeof(StrandBlock_t), NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(StrandBlock_t), NULL);
380      CheckErrorCode (sts, "P2 error Setting OpenCL local strand blockargument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL local strand blockargument\n");
     *argCount+=1;  
381    
382      sts = clSetKernelArg (kernel,  *(argCount), sizeof(int) * blk_sz, NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(int) * blk_sz, NULL);
383      CheckErrorCode (sts, "P2 error Setting OpenCL local preStable argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL local preStable argument\n");
     *argCount+=1;  
384    
385      sts = clSetKernelArg (kernel,  *(argCount), sizeof(int) * blk_sz, NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(int) * blk_sz, NULL);
386      CheckErrorCode (sts, "P2 error Setting OpenCL local preDead argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL local preDead argument\n");
     *argCount+=1;  
387    
388      sts = clSetKernelArg (kernel,  *(argCount), sizeof(int) , NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(int) , NULL);
389      CheckErrorCode (sts, "P2 error Setting OpenCL local numActive argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL local numActive argument\n");
     *argCount+=1;  
   
390    
391      sts = clSetKernelArg (kernel,  *(argCount), sizeof(int) * blk_sz , NULL);      sts = clSetKernelArg (kernel,  count++, sizeof(int) * blk_sz , NULL);
392      CheckErrorCode (sts, "P2 error Setting OpenCL local temporary array for the prefix scan of preStable preDead argument\n");      CheckErrorCode (sts, "P2 error Setting OpenCL local temporary array for the prefix scan of preStable preDead argument\n");
     *argCount+=1;  
393    
394        *argCount = count;
395  }  }
396    
397  static void CheckErrorCode (cl_int sts, const char *msg)  static void CheckErrorCode (cl_int sts, const char *msg)
# Line 520  Line 487 
487      for (i = 0;  i < clInfo->platforms[0].numDevices;  i++) {      for (i = 0;  i < clInfo->platforms[0].numDevices;  i++) {
488          if (isGPUDevice (&(clInfo->platforms[0].devices[i]))          if (isGPUDevice (&(clInfo->platforms[0].devices[i]))
489          &&  clInfo->platforms[0].devices[i].isAvail) {          &&  clInfo->platforms[0].devices[i].isAvail) {
   
490              dev = &(clInfo->platforms[0].devices[i]);              dev = &(clInfo->platforms[0].devices[i]);
491              clInfo->mainDeviceIdx = i;              clInfo->mainDeviceIdx = i;
492              break;              break;
# Line 551  Line 517 
517      char *updateSource = LoadSource (fname);      char *updateSource = LoadSource (fname);
518      free (fname);      free (fname);
519    
520      int kLen = strlen(DIDEROT_INCLUDE_PATH) + strlen(kernelsFile) + 2;    // load scheduler kernels
521      char *kname = (char *)CheckedAlloc(kLen);      char *kernelsSource = LoadSource(KERNELS_FILE_PATH);
     snprintf(kname, kLen, "%s/%s",DIDEROT_INCLUDE_PATH,kernelsFile);  
     char *kernelsSource = LoadSource(kname);  
     free(kname);  
522    
523      const char *src[2] = {updateSource, kernelsSource};      const char *src[2] = {updateSource, kernelsSource};
524      cl_program prog = clCreateProgramWithSource(cxt, 2, src, NULL, &sts);      cl_program prog = clCreateProgramWithSource(cxt, 2, src, NULL, &sts);
525        free (updateSource);
526        free (kernelsSource);
527      if (sts != CL_SUCCESS) {      if (sts != CL_SUCCESS) {
528          fprintf (stderr, "error creating program\n");          fprintf (stderr, "error creating program\n");
529          return false;          return false;

Legend:
Removed from v.1491  
changed lines
  Added in v.1492

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