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 1461, Mon Aug 8 04:50:12 2011 UTC revision 1462, Tue Aug 9 07:22:45 2011 UTC
# Line 108  Line 108 
108    
109      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
110    
     // hack to make the invariant part of the state the same in both copies  
     memcpy (wrld->outState, wrld->inState, wrld->strandSize);  
111    
112      if (! InitCL(clInfo, wrld))      if (! InitCL(clInfo, wrld))
113          exit (1);          exit (1);
114    
115    
116        /* Conversion of strands from their host types to their shadow types */
117        void * shadowInState = malloc(Diderot_Strands[0]->shadowStrandSzb * wrld->numStrands);
118        uint8_t *strandPtr = (uint8_t *)wrld->inState;
119        uint8_t *strandShadowPtr = (uint8_t *)shadowInState;
120        size_t shadowSize = Diderot_Strands[0]->shadowStrandSzb * wrld->numStrands;
121    
122        for (int i = 0;  i < wrld->numStrands;  i++, strandPtr += Diderot_Strands[0]->stateSzb,strandShadowPtr+= Diderot_Strands[0]->shadowStrandSzb) {
123                Diderot_Strands[0]->strandCopy (strandPtr,strandShadowPtr);
124        }
125        free((uint8_t *)wrld->inState);
126        wrld->inState = shadowInState;
127    
128    
129    
130        // hack to make the invariant part of the state the same in both copies
131        memcpy (wrld->outState, wrld->inState, shadowSize);
132    
133    
134    
135      int argCount = 0;      int argCount = 0;
136      cl_int sts = CL_SUCCESS;      cl_int sts = CL_SUCCESS;
137      UpdateKernel_Args updateArgs;      UpdateKernel_Args updateArgs;
# Line 121  Line 139 
139    
140    
141      updateArgs.stateOutMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,      updateArgs.stateOutMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
142          wrld->strandSize, wrld->outState, &sts);          shadowSize, wrld->outState, &sts);
143      checkErrorCode(sts,"error creating OpenCL strand in-state buffer\n");      checkErrorCode(sts,"error creating OpenCL strand in-state buffer\n");
144    
145    
# Line 138  Line 156 
156      int workGroupSize = 0;      int workGroupSize = 0;
157      int * workQueue;      int * workQueue;
158    
159    
160   if (wrld->nDims == 1) {   if (wrld->nDims == 1) {
161          globalWorkSize[0] = wrld->size[0] + (localWorkSize[0] - (wrld->size[0] & (localWorkSize[0] - 1)));          globalWorkSize[0] = wrld->size[0] + (localWorkSize[0] - (wrld->size[0] & (localWorkSize[0] - 1)));
162          numberOfWorkGroups = (globalWorkSize[0] / localWorkSize[0]);          numberOfWorkGroups = (globalWorkSize[0] / localWorkSize[0]);
# Line 150  Line 169 
169          workGroupSize = localWorkSize[0] * localWorkSize[1];          workGroupSize = localWorkSize[0] * localWorkSize[1];
170      }      }
171      else if (wrld->nDims == 3) {      else if (wrld->nDims == 3) {
172            localWorkSize[0] = 8;
173            localWorkSize[1] = 8;
174            localWorkSize[2] = 8;
175          globalWorkSize[0] = wrld->size[0] + (localWorkSize[0] - (wrld->size[0] & (localWorkSize[0] - 1)));          globalWorkSize[0] = wrld->size[0] + (localWorkSize[0] - (wrld->size[0] & (localWorkSize[0] - 1)));
176          globalWorkSize[1] = wrld->size[1] + (localWorkSize[1] - (wrld->size[1] & (localWorkSize[1] - 1)));          globalWorkSize[1] = wrld->size[1] + (localWorkSize[1] - (wrld->size[1] & (localWorkSize[1] - 1)));
177          globalWorkSize[2] = wrld->size[2] + (localWorkSize[2] - (wrld->size[2] & (localWorkSize[2] - 1)));          globalWorkSize[2] = wrld->size[2] + (localWorkSize[2] - (wrld->size[2] & (localWorkSize[2] - 1)));
# Line 160  Line 182 
182          assert (0);          assert (0);
183      }      }
184    
185    
186      updateArgs.numActiveMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,      updateArgs.numActiveMem = clCreateBuffer (wrld->context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
187       sizeof(int), &numActive, &sts);       sizeof(int), &numActive, &sts);
188       checkErrorCode(sts,"error creating OpenCL number of avaliable  buffer\n");       checkErrorCode(sts,"error creating OpenCL number of avaliable  buffer\n");
# Line 198  Line 221 
221    
222       clFinish(wrld->cmdQ);       clFinish(wrld->cmdQ);
223    
224    
225      while(numActive > 0)      while(numActive > 0)
226      {      {
227          sts = clEnqueueTask(wrld->cmdQ, wrld->kernel[0], 0, NULL, NULL);          sts = clEnqueueTask(wrld->cmdQ, wrld->kernel[0], 0, NULL, NULL);
# Line 207  Line 231 
231    
232          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1], wrld->nDims, NULL, globalWorkSize,          sts = clEnqueueNDRangeKernel(wrld->cmdQ, wrld->kernel[1], wrld->nDims, NULL, globalWorkSize,
233                localWorkSize, 0, NULL, NULL);                localWorkSize, 0, NULL, NULL);
234    
235          checkErrorCode(sts, "error in executing update kernel\n");          checkErrorCode(sts, "error in executing update kernel\n");
236          clFinish(wrld->cmdQ);          clFinish(wrld->cmdQ);
237    
238          sts = clEnqueueReadBuffer(wrld->cmdQ, updateArgs.numActiveMem, CL_TRUE, 0, sizeof(int),          sts = clEnqueueReadBuffer(wrld->cmdQ, updateArgs.numActiveMem, CL_TRUE, 0, sizeof(int),
239             &numActive, 0, NULL, NULL);             &numActive, 0, NULL, NULL);
240          checkErrorCode(sts, "error in reading back number of active strands");  
241            checkErrorCode(sts, "error in reading back number of active strands: %d");
242    
243          clFinish(wrld->cmdQ);          clFinish(wrld->cmdQ);
244    
245      }      }
246    
247      sts = clEnqueueReadBuffer(wrld->cmdQ, updateArgs.stateOutMem, CL_TRUE, 0, wrld->strandSize,      sts = clEnqueueReadBuffer(wrld->cmdQ, updateArgs.stateOutMem, CL_TRUE, 0,shadowSize,
248              wrld->outState, 0, NULL, NULL);              wrld->outState, 0, NULL, NULL);
249      checkErrorCode(sts, "error in reading back output\n");      checkErrorCode(sts, "error in reading back output\n");
250    
# Line 250  Line 276 
276          exit (8);          exit (8);
277      }      }
278    
279      uint8_t *strandPtr = (uint8_t *)wrld->outState;      strandPtr = (uint8_t *)wrld->outState;
280      for (int i = 0;  i < wrld->numStrands;  i++, strandPtr += Diderot_Strands[0]->stateSzb) {      for (int i = 0;  i < wrld->numStrands;  i++, strandPtr += Diderot_Strands[0]->shadowStrandSzb) {
281          if (wrld->status[i] == DIDEROT_STABLE)          if (wrld->status[i] == DIDEROT_STABLE)
282              Diderot_Strands[0]->print (outS, strandPtr);              Diderot_Strands[0]->print (outS, strandPtr);
283      }      }
# Line 497  Line 523 
523    /*  wrld->inState = (void **) malloc (numStrands * sizeof(void *));    /*  wrld->inState = (void **) malloc (numStrands * sizeof(void *));
524      wrld->outState = (void **) malloc (numStrands * sizeof(void *)); */      wrld->outState = (void **) malloc (numStrands * sizeof(void *)); */
525      wrld->inState =  malloc (wrld->strandSize);      wrld->inState =  malloc (wrld->strandSize);
526      wrld->outState = malloc (wrld->strandSize);      wrld->outState = malloc (strand->shadowStrandSzb * numStrands);
527      wrld->status = (int *) malloc (numStrands * sizeof(int));      wrld->status = (int *) malloc (numStrands * sizeof(int));
528      if ((wrld->inState == 0) || (wrld->outState == 0) || (wrld->status == 0)) {      if ((wrld->inState == 0) || (wrld->outState == 0) || (wrld->status == 0)) {
529          fprintf (stderr, "unable to allocate strand states\n");          fprintf (stderr, "unable to allocate strand states\n");

Legend:
Removed from v.1461  
changed lines
  Added in v.1462

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