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

SCM Repository

[diderot] Diff of /branches/pure-cfg/src/include/Diderot/cl-scheduler.h
ViewVC logotype

Diff of /branches/pure-cfg/src/include/Diderot/cl-scheduler.h

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 51  Line 51 
51    
52      //Check to see if the index retrieve is less than the      //Check to see if the index retrieve is less than the
53      //number of blocks on the queue      //number of blocks on the queue
54      if(blkIdx < sched->queueSize)      if (blkIdx < sched->queueSize) {
     {  
55          //Copy over the global data for the block into local memory          //Copy over the global data for the block into local memory
56          lBlk->blkIdx = (blocks[queue[blkIdx]]).blkIdx;          lBlk->blkIdx = (blocks[queue[blkIdx]]).blkIdx;
57          lBlk->nActive = (blocks[queue[blkIdx]]).nActive;          lBlk->nActive = (blocks[queue[blkIdx]]).nActive;
# Line 62  Line 61 
61    
62          //If the his block has no active strands then check to see if          //If the his block has no active strands then check to see if
63          // there are strands still waiting to be processed          // there are strands still waiting to be processed
64          if(lBlk->nActive == 0)          if (lBlk->nActive == 0)  {
         {  
65              //Get a block of strands              //Get a block of strands
66              int startStrandIdx = atom_add(&(sched->nextStrand[0]), BLK_SZ);              int startStrandIdx = atom_add(&(sched->nextStrand[0]), BLK_SZ);
67              int idx = 0, count = 0;              int idx = 0, count = 0;
68              //For each work-item in the workgroup              //For each work-item in the workgroup
69              for(int i = startStrandIdx, count = 0; i < sched->numStrands && count < BLK_SZ; i++,count++)              for (int i = startStrandIdx, count = 0; i < sched->numStrands && count < BLK_SZ; i++,count++)  {
             {  
70                  //Increment the number of active strands and assign each work-item a strand to process                  //Increment the number of active strands and assign each work-item a strand to process
71                  lBlk->nActive++;                  lBlk->nActive++;
72                  blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;                  blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;
# Line 80  Line 77 
77                  lBlk->nActive = -1;                  lBlk->nActive = -1;
78          }          }
79      }      }
80      //IF there isn't any queue items then notify the the workgroup      else // if there isn't any queue items then notify the the workgroup
     else  
81          lBlk->nActive = -1;          lBlk->nActive = -1;
82  }  }
83    
# Line 101  Line 97 
97      int blkIdx = atom_inc(&(sched->sId));      int blkIdx = atom_inc(&(sched->sId));
98    
99      //Check to make sure the idx retrieve is a valid for the list      //Check to make sure the idx retrieve is a valid for the list
100      if(blkIdx < sched->todoSize)      if (blkIdx < sched->todoSize) {
     {  
101          //Copy over the global data into the local data.          //Copy over the global data into the local data.
102          lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;          lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;
103          lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;          lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;
# Line 114  Line 109 
109          lBlk->nActive = -1;          lBlk->nActive = -1;
110  }  }
111    
112  // Exclusive Scan that requires the number of work-items to be less than // the maximum number of work-items for a single work-group And the only // one work group to be executed. inline void scan ( __global int * input,  // Exclusive Scan that requires the number of work-items to be less than
113                __global int * status,              __local int * output,  // the maximum number of work-items for a single work-group and the only
114                StrandStatus_t statusTy,              int n,  // one work group to be executed.
115                __local int * temp) {    int thid = get_local_id(0);    int pout = 0, pin = 1;  inline void scan (
116        //load input into local memory    // Exclusive scan: shift right by one and set first element to 0      __global int        *input,
117        __global int        *status,
118        __local int         *output,
119        StrandStatus_t      statusTy,
120        int                 n,
121        __local int         *temp)
122    {
123        int thid = get_local_id(0);
124        int pout = 0, pin = 1;
125    
126      // load input into local memory
127      // Exclusive scan: shift right by one and set first element to 0
128      if(thid == 0 || status[(*(input + thid))] != statusTy)      if(thid == 0 || status[(*(input + thid))] != statusTy)
129        temp[thid] = 0;        temp[thid] = 0;
130      else      temp[thid] = 1;      else
131            barrier(CLK_LOCAL_MEM_FENCE);    for(int offset = 1; offset < n; offset *= 2)    {          pout = 1 - pout;        pin = 1 - pout;                if(thid >= offset)            temp[pout * n + thid] = temp[pin * n + thid] + temp[pin * n + thid - offset];        else            temp[pout * n + thid] = temp[pin * n + thid];            barrier(CLK_LOCAL_MEM_FENCE);    }    output[thid] = temp[pout* n + thid]; }          temp[thid] = 1;
132    
133        barrier(CLK_LOCAL_MEM_FENCE);
134    
135        for (int offset = 1; offset < n; offset *= 2) {
136            pout = 1 - pout;
137            pin = 1 - pout;
138    
139            if(thid >= offset)
140                temp[pout * n + thid] = temp[pin * n + thid] + temp[pin * n + thid - offset];
141            else
142                temp[pout * n + thid] = temp[pin * n + thid];
143    
144            barrier(CLK_LOCAL_MEM_FENCE);
145        }
146        output[thid] = temp[pout* n + thid];
147    }
148  #endif /* !_DIDEROT_CL_SCHEDULER_H_ */  #endif /* !_DIDEROT_CL_SCHEDULER_H_ */

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