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

SCM Repository

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

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

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

revision 1678, Sun Dec 18 13:17:09 2011 UTC revision 1679, Mon Dec 19 20:31:04 2011 UTC
# Line 155  Line 155 
155      if(sched->clearQueueSz == 1) {      if(sched->clearQueueSz == 1) {
156          sched->clearQueueSz = 0;          sched->clearQueueSz = 0;
157          sched->queueSize = 0;          sched->queueSize = 0;
158      } else {      }
159        else {
160          sched->clearQueueSz = 1;          sched->clearQueueSz = 1;
161          sched->todoSize = 0;          sched->todoSize = 0;
162      }      }
# Line 189  Line 190 
190      int qIdx;      int qIdx;
191    
192      do {      do {
193          if (get_local_id(0) == 0) {          if (id == 0) {
194             bool done;             bool done;
195             do {             do {
196              GetBlockFromTodoList(sched,todoList,blocks,bp);              GetBlockFromTodoList(sched,todoList,blocks,bp);
# Line 207  Line 208 
208                     done = false;                     done = false;
209                  }                  }
210                }                }
211                } while (!done);
             }  
             while (!done);  
212          }          }
213          barrier (CLK_LOCAL_MEM_FENCE);          barrier (CLK_LOCAL_MEM_FENCE);
214    
# Line 224  Line 223 
223            // compaction            // compaction
224            // with these arrays, we can then compute the new index of each slot in            // with these arrays, we can then compute the new index of each slot in
225            // parallel as follows            // parallel as follows
226                int newIdx;
227              switch (status[idx]) {              switch (status[idx]) {
228                case DIDEROT_ACTIVE:                case DIDEROT_ACTIVE:
229                  id = id - preStable[id] - preDead[id];                  newIdx = id - preStable[id] - preDead[id];
230                  break;                  break;
231                case DIDEROT_STABLE:                case DIDEROT_STABLE:
232                  id = bp->nActive + preStable[id];                  newIdx = bp->nActive + preStable[id];
233                  break;                  break;
234                case DIDEROT_DIE:                case DIDEROT_DIE:
235                  id = bp->nActive + preStable[id] + preDead[id];                  newIdx = bp->nActive + preStable[id] + preDead[id];
236                  break;                  break;
237              }              }
238              blockIndxs[bp->blkIdx * BLK_SZ + id] = idx;              blockIndxs[bp->blkIdx * BLK_SZ + newIdx] = idx;
239    
240              barrier (CLK_LOCAL_MEM_FENCE);              barrier (CLK_LOCAL_MEM_FENCE);
241              if (get_local_id(0) == 0) {  // note that id ?= get_local_id(0)              if (id == 0) {
242                  // put bp on scheduling queue                  // put bp on scheduling queue
243                 qIdx = atom_inc(&(sched->queueSize));                 qIdx = atom_inc(&(sched->queueSize));
244                 bp->nStabilizing = 0;                 bp->nStabilizing = 0;
Legend:
Removed from v.1678  
changed lines
  Added in v.1679
root@smlnj-gforge.cs.uchicago.edu
 ViewVC Help
Powered by ViewVC 1.0.0  
Powered By FusionForge