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 1518, Sat Oct 8 02:39:41 2011 UTC revision 1519, Sun Oct 9 18:30:16 2011 UTC
# Line 10  Line 10 
10    
11  #include <Diderot/diderot.h>  #include <Diderot/diderot.h>
12  #include "clinfo.h"  #include "clinfo.h"
 #include <string.h>  
 #include <stdio.h>  
 #include <assert.h>  
13  #include <sys/sysctl.h>  #include <sys/sysctl.h>
14  #include <sys/stat.h>  #include <sys/stat.h>
15    
# Line 53  Line 50 
50  } GPUKernel_t;  } GPUKernel_t;
51    
52  struct struct_world {  struct struct_world {
53      const char          *name;          //!< the program name      STRUCT_WORLD_PREFIX
54      bool                isArray;        //!< is the initialization an array or collection?  // FIXME: document these fields!
     uint32_t            nDims;          //!< depth of iteration nesting  
     int32_t             *base;          //!< nDims array of base indices  
     uint32_t            *size;          //!< nDims array of iteration sizes  
     uint32_t            numStrands;     //!< number of strands in the world  
     uint32_t            strandSize;     //!< the sizeof of the strand buffers  
55      void                *inState;      void                *inState;
56      void                *outState;      void                *outState;
     uint32_t            *status;        //!< array of strand status flags  
57      cl_device_id        device;         //!< OpenCL device      cl_device_id        device;         //!< OpenCL device
58      cl_context          context;        //!< OpenCL execution context      cl_context          context;        //!< OpenCL execution context
59      cl_command_queue    cmdQ;           //!< OpenCL command queue      cl_command_queue    cmdQ;           //!< OpenCL command queue
60      GPUKernel_t         kernel[3];      //!< OpenCL kernel that implements the program      GPUKernel_t         kernel[3];      //!< OpenCL kernel that implements the program
61  };  };
62    
63    // FIXME: need documentation for this stuct!
64  typedef struct {  typedef struct {
65     cl_mem schedMem;     cl_mem schedMem;
66     cl_mem outMem;     cl_mem outMem;
# Line 78  Line 70 
70     cl_mem queueMem;     cl_mem queueMem;
71     cl_mem todoMem;     cl_mem todoMem;
72     cl_mem statusMem;     cl_mem statusMem;
73  } Kernel_Args_t;  } KernelArgs_t;
74    
75  static bool InitCL (CLInfo_t *clInfo, Diderot_World_t *wrld);  static bool InitCL (CLInfo_t *clInfo, Diderot_World_t *wrld);
76  static void CheckErrorCode (cl_int sts, const char * msg);  static void CheckErrorCode (cl_int sts, const char * msg);
77  static void SetPhase1Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args);  static void SetPhase1Args (cl_kernel kernel, int *argCount, KernelArgs_t *args);
78  static void SetPhase2Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args, int blk_sz);  static void SetPhase2Args (cl_kernel kernel, int *argCount, KernelArgs_t *args, int blk_sz);
79  static void SetScheduleKernelArgs(cl_kernel kernel, int *argCount, Kernel_Args_t *args);  static void SetScheduleKernelArgs (cl_kernel kernel, int *argCount, KernelArgs_t *args);
80    
81  static bool CreateKernel (cl_device_id dev, cl_program prog, const char *name, GPUKernel_t *kern);  static bool CreateKernel (cl_device_id dev, cl_program prog, const char *name, GPUKernel_t *kern);
82    
# Line 136  Line 128 
128          strandShadowPtr += Diderot_Strands[0]->shadowStrandSzb;          strandShadowPtr += Diderot_Strands[0]->shadowStrandSzb;
129          strandShadowOutPtr += Diderot_Strands[0]->shadowStrandSzb;          strandShadowOutPtr += Diderot_Strands[0]->shadowStrandSzb;
130      }      }
131    // FIXME: it is confusing to use the inState/outState pointers for two different purposes.
132      free (wrld->inState);      free (wrld->inState);
133      free(wrld->outState);      free(wrld->outState);
134      wrld->inState = shadowInState;      wrld->inState = shadowInState;
# Line 147  Line 140 
140      int argCount = 0;      int argCount = 0;
141      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
142    
143      Kernel_Args_t kernelArgs;      KernelArgs_t kernelArgs;
144      SchedState_t scheduler;      SchedState_t scheduler;
145    
146      scheduler.numStrands = wrld->numStrands;      scheduler.numStrands = wrld->numStrands;
# Line 235  Line 228 
228      clFinish(wrld->cmdQ);      clFinish(wrld->cmdQ);
229    
230      while (scheduler.numAvailable > 0) {      while (scheduler.numAvailable > 0) {
   
231        //Runs the update kernel on all  strands        //Runs the update kernel on all  strands
232          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[0].kern, 1, NULL,          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[0].kern, 1, NULL,
233              globalWorkSize, localWorkSize, 0, NULL, NULL);              globalWorkSize, localWorkSize, 0, NULL, NULL);
# Line 244  Line 236 
236         sts = clEnqueueTask(wrld->cmdQ,wrld->kernel[2].kern,0,NULL,NULL);         sts = clEnqueueTask(wrld->cmdQ,wrld->kernel[2].kern,0,NULL,NULL);
237         CheckErrorCode (sts, "error in executing scheduler update kernel before compaction\n");         CheckErrorCode (sts, "error in executing scheduler update kernel before compaction\n");
238    
   
239         // Run the compaction kernel on all strands         // Run the compaction kernel on all strands
240         sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1].kern, 1, NULL,         sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1].kern, 1, NULL,
241                  globalWorkSize, localWorkSize, 0, NULL, NULL);                  globalWorkSize, localWorkSize, 0, NULL, NULL);
242         CheckErrorCode (sts, "error ccccn executing compaction kernel\n");         CheckErrorCode (sts, "error ccccn executing compaction kernel\n");
243    
   
244         sts = clEnqueueTask(wrld->cmdQ,wrld->kernel[2].kern,0,NULL,NULL);         sts = clEnqueueTask(wrld->cmdQ,wrld->kernel[2].kern,0,NULL,NULL);
245         CheckErrorCode (sts, "error in executing scheduler update kernel after compaction\n");         CheckErrorCode (sts, "error in executing scheduler update kernel after compaction\n");
246    
   
247         sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),         sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.schedMem, CL_TRUE, 0, sizeof(SchedState_t),
248                  &scheduler, 0, NULL, NULL);                  &scheduler, 0, NULL, NULL);
   
249         CheckErrorCode (sts, "error reading back scheduler information\n");         CheckErrorCode (sts, "error reading back scheduler information\n");
   
   
250      }      }
251    
252      sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.outMem, CL_TRUE, 0,shadowSize,      sts = clEnqueueReadBuffer(wrld->cmdQ, kernelArgs.outMem, CL_TRUE, 0,shadowSize,
# Line 288  Line 274 
274    
275    // here we have the final state of all of the strands in the "in" buffer    // here we have the final state of all of the strands in the "in" buffer
276    // output the final strand states    // output the final strand states
     Output_Args_t outArgs;  
     outArgs.name = wrld->name;  
     outArgs.isArray = wrld->isArray;  
     outArgs.numStrands = wrld->numStrands;  
     outArgs.status = wrld->status;  
     outArgs.inState = wrld->outState;  
     outArgs.nDims = wrld->nDims;  
     outArgs.size = wrld->size;  
     outArgs.outputSzb = Diderot_Strands[0]->shadowStrandSzb;  
277      if (NrrdOutputFlg)      if (NrrdOutputFlg)
278          Diderot_Output (&outArgs);          Diderot_Output (wrld, Diderot_Strands[0]->shadowStrandSzb);
279      else      else
280          Diderot_Print (&outArgs);          Diderot_Print (wrld);
281    
282      Diderot_Shutdown (wrld);      Diderot_Shutdown (wrld);
283    
# Line 308  Line 285 
285    
286  }  }
287    
288  static void SetPhase1Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args)  static void SetPhase1Args (cl_kernel kernel, int *argCount, KernelArgs_t *args)
289  {  {
290      int count = *argCount;      int count = *argCount;
291      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
# Line 344  Line 321 
321    
322  }  }
323    
324  static void SetPhase2Args (cl_kernel kernel, int *argCount, Kernel_Args_t *args, int blk_sz)  static void SetPhase2Args (cl_kernel kernel, int *argCount, KernelArgs_t *args, int blk_sz)
325  {  {
326      int count = *argCount;      int count = *argCount;
327      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
# Line 382  Line 359 
359    
360      *argCount = count;      *argCount = count;
361  }  }
362  static void SetScheduleKernelArgs(cl_kernel kernel, int *argCount, Kernel_Args_t *args)  static void SetScheduleKernelArgs(cl_kernel kernel, int *argCount, KernelArgs_t *args)
363  {  {
364      int count = *argCount;      int count = *argCount;
365      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
# Line 612  Line 589 
589    
590    // allocate the strand state pointers    // allocate the strand state pointers
591      wrld->numStrands = numStrands;      wrld->numStrands = numStrands;
592      wrld->strandSize = strand->stateSzb * numStrands;      wrld->inState = CheckedAlloc (strand->stateSzb * numStrands);
 /*  
     wrld->inState = NEWVEC(void *, numStrands);  
     wrld->outState = NEWVEC(void *, numStrands);  
 */  
     wrld->inState =  CheckedAlloc (wrld->strandSize);  
593      wrld->outState = CheckedAlloc (strand->shadowStrandSzb * numStrands);      wrld->outState = CheckedAlloc (strand->shadowStrandSzb * numStrands);
594      wrld->status = NEWVEC(uint32_t, numStrands);      wrld->status = NEWVEC(StatusInt_t, numStrands);
     if ((wrld->inState == 0) || (wrld->outState == 0) || (wrld->status == 0)) {  
         fprintf (stderr, "unable to allocate strand states\n");  
         exit (1);  
     }  
595    
596    // initialize strand state pointers etc.    // initialize strand state pointers etc.
597      for (size_t i = 0;  i < numStrands;  i++) {      for (size_t i = 0;  i < numStrands;  i++) {
        // wrld->inState[i] = Diderot_AllocStrand (strand);  
        // wrld->outState[i] = Diderot_AllocStrand (strand);  
598          wrld->status[i] = DIDEROT_ACTIVE;          wrld->status[i] = DIDEROT_ACTIVE;
599      }      }
600    
# Line 640  Line 606 
606  void *Diderot_InState (Diderot_World_t *wrld, uint32_t i)  void *Diderot_InState (Diderot_World_t *wrld, uint32_t i)
607  {  {
608      assert (i < wrld->numStrands);      assert (i < wrld->numStrands);
609      return wrld->inState + i;      return &(wrld->inState[i]);
610  }  }
611    
612  void *Diderot_OutState (Diderot_World_t *wrld, uint32_t i)  void *Diderot_OutState (Diderot_World_t *wrld, uint32_t i)
613  {  {
614      assert (i < wrld->numStrands);      assert (i < wrld->numStrands);
615      return &wrld->outState[i];      return &(wrld->outState[i]);
 }  
   
 bool Diderot_IsActive (Diderot_World_t *wrld, uint32_t i)  
 {  
     assert (i < wrld->numStrands);  
     return !wrld->status[i];  
616  }  }
617    
618  /***** Support for shadow image values *****/  /***** Support for shadow image values *****/

Legend:
Removed from v.1518  
changed lines
  Added in v.1519

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