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

SCM Repository

[diderot] Diff of /branches/staging/src/compiler/cl-target/fragments/sched.in
ViewVC logotype

Diff of /branches/staging/src/compiler/cl-target/fragments/sched.in

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

revision 2005, Fri Oct 5 11:57:00 2012 UTC revision 2006, Fri Oct 5 11:57:14 2012 UTC
# Line 147  Line 147 
147      output[thid] = temp[pout* n + thid];      output[thid] = temp[pout* n + thid];
148  }  }
149    
150  // Scheduler Meta-Clearing Kenerl: Clears the queue sizes and strand id acculator  // Scheduler Meta-Clearing Kernel: Clears the queue sizes and strand id acculator
151  //! \param sched global scheduler state  //! \param sched global scheduler state
152  __kernel void Diderot_SchedUpdateKernel (__global SchedState_t *sched)  __kernel void Diderot_SchedUpdateKernel (__global SchedState_t *sched)
153  {  {
# Line 183  Line 183 
183      __local StrandBlock_t *bp,      __local StrandBlock_t *bp,
184      __local int *preStable,      __local int *preStable,
185      __local int *preDead,      __local int *preDead,
186      __local int *prefixScanTemp      __local int *prefixScanTemp)
     )  
187  {  {
188      int id = get_local_id(0);      int id = get_local_id(0);
189      int qIdx;      int qIdx;
190    
191      do {      do {
192          if (id == 0) {          if (id == 0) {
            bool done;  
            do {  
193              GetBlockFromTodoList(sched,todoList,blocks,bp);              GetBlockFromTodoList(sched,todoList,blocks,bp);
194              done = true;            // bp->nAvail < 0 if there are no more blocks to process
195              if (bp->nActive > 0) {              if (bp->nActive > 0) {
196                  bp->nActive =  bp->nActive - (bp->nStabilizing + bp->nDying);                  bp->nActive =  bp->nActive - (bp->nStabilizing + bp->nDying);
197                  bp->nDead += bp->nDying;                  bp->nDead += bp->nDying;
                 if (bp->nActive == 0) {  
                    //Check to see if this block can execute more strands for the  
                    //next iteration  
                    GrabWork(sched,blockIndxs,bp);  
                    bp->nStabilizing = 0;  
                    bp->nDying = 0;  
                    StrandBlock_Copy(&blocks[bp->blkIdx],bp);  
                    done = false;  
                 }  
198                }                }
             } while (!done);  
199          }          }
200          barrier (CLK_LOCAL_MEM_FENCE);          barrier (CLK_LOCAL_MEM_FENCE);
201    
# Line 247  Line 234 
234                  queue[qIdx] = bp->blkIdx;                  queue[qIdx] = bp->blkIdx;
235              }              }
236          }          }
237      } while (bp->nActive > 0);      } while (bp->nActive >= 0);
238    
239  }  }

Legend:
Removed from v.2005  
changed lines
  Added in v.2006

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