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 1501, Tue Sep 13 19:33:15 2011 UTC revision 1513, Fri Oct 7 14:54:06 2011 UTC
# Line 15  Line 15 
15  #  include "cl-types.h"  #  include "cl-types.h"
16  #endif  #endif
17    
 /* Scheduler meta-information */  
 typedef struct {  
     int    blkIdx;         // the id of this block  
     int    nActive;        // number of active (status != DIE or STABLE) strands  
     int    nDead;          // number of strands in the DIE state  
     int    nStabilizing;   // number of new strands in the STABILIZE state  
     int    nDying;         // number of new strands in the DIE state  
 } StrandBlock_t;  
   
 typedef struct {  
     int numStrands;                 // number of strands  
     int sId;                        // the index accumlator for the todo list or queue  
     int nextStrand;              // index of the next strand to retrieve from the pool  
     int clearQueueSz;              // an indicator on whether the queue size should be cleared  
     int queueSize;                  // number of blocks on the scheduler's queue  
     int todoSize;                   // number of blocks on the scheduler's todo list  
     int numAvailable;            // number of active strands left to process  
 } SchedState_t;  
   
   
 //! \brief Get a block of strands from the scheduler pool.  
 //! \param sched global scheduler state  
 //! \param blocks global array of strand blocks  
 //! \param lBlk pointer to global storage for the strand block  
 //! \return lBlk if there is a block to be scheduled, and 0 otherwise  
 inline void GetBlock (  
     __global SchedState_t       *sched,  
     __global int                *blockIndxs,  
     __global int                *queue,  
     __global StrandBlock_t      *blocks,  
     __local StrandBlock_t       *lBlk)  
 {  
   // Retrieve a block from the list  
     int blkIdx = atom_inc(&(sched->sId));  
   
   // Check to see if the index retrieve is less than the  
   // number of blocks on the queue  
     if (blkIdx < sched->queueSize) {  
       // Copy over the global data for the block into local memory  
         lBlk->blkIdx = (blocks[queue[blkIdx]]).blkIdx;  
         lBlk->nActive = (blocks[queue[blkIdx]]).nActive;  
         lBlk->nDead = (blocks[queue[blkIdx]]).nDead;  
         lBlk->nStabilizing = (blocks[queue[blkIdx]]).nStabilizing;  
         lBlk->nDying = (blocks[queue[blkIdx]]).nDying;  
   
       // If the his block has no active strands then check to see if  
       // there are strands still waiting to be processed  
         if (lBlk->nActive == 0)  {  
           // Get a block of strands  
             int startStrandIdx = atom_add(&sched->nextStrand, BLK_SZ);  
             int idx = 0, count = 0;  
           // For each work-item in the workgroup  
             for (int i = startStrandIdx, count = 0; i < sched->numStrands && count < BLK_SZ; i++,count++)  {  
               // Increment the number of active strands and assign each work-item a strand to process  
                 lBlk->nActive++;  
                 blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;  
                 idx++;  
             }  
           // If there are no strands left to process than make this block inactive  
             if(idx == 0)  
                 lBlk->nActive = -1;  
         }  
     }  
     else // if there isn't any queue items then notify the the workgroup  
         lBlk->nActive = -1;  
 }  
 //! \brief Copies the strand black from __local space to __global space  
 //! \param sched global scheduler state  
 //! \param blocks global array of strand blocks  
 //! \param lBlk pointer to global storage for the strand block  
 //! \return lBlk if there is a block to be scheduled, and 0 otherwise  
 inline void StrandBlock_Copy (__global StrandBlock_t * gBlock,  
                                   __local StrandBlock_t * lBlock)  
 {  
    gBlock->nActive = lBlock->nActive;  
    gBlock->nDead = lBlock->nDead;  
    gBlock->nStabilizing = lBlock->nStabilizing;  
    gBlock->nDying = lBlock->nDying;  
 }  
 //! \brief Get a block of strands from the Todolist scheduler pool.  
 //! \param sched global scheduler state  
 //! \param blocks global array of strand blocks  
 //! \param lBlk pointer to global storage for the strand block  
 //! \return lBlk if there is a block to be scheduled, and 0 otherwise  
 inline void GetBlockFromTodoList (  
     __global SchedState_t *sched,  
     __global int *blockIndxs,  
     __global int *todoList,  
     __global StrandBlock_t *blocks,  
     __local StrandBlock_t *lBlk)  
 {  
   // Retrieve a block from the todolist  
     int blkIdx = atom_inc(&(sched->sId));  
   
   // Check to make sure the idx retrieve is a valid for the list  
     if (blkIdx < sched->todoSize) {  
       // Copy over the global data into the local data.  
         lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;  
         lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;  
         lBlk->nDead = (blocks[todoList[blkIdx]]).nDead;  
         lBlk->nStabilizing = (blocks[todoList[blkIdx]]).nStabilizing;  
         lBlk->nDying = (blocks[todoList[blkIdx]]).nDying;  
     }  
     else  
         lBlk->nActive = -1;  
 }  
   
 // 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,  
     __global int        *status,  
     __local int         *output,  
     StrandStatus_t      statusTy,  
     int                 n,  
     __local int         *temp)  
 {  
     int thid = get_local_id(0);  
     int pout = 0, pin = 1;  
   
   // load input into local memory  
   // Exclusive scan: shift right by one and set first element to 0  
     if (thid == 0 || status[(*(input + thid - 1))] != statusTy)  
         temp[thid] = 0;  
     else  
         temp[thid] = 1;  
   
     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];  
 }  
18  #endif /* !_DIDEROT_CL_SCHEDULER_H_ */  #endif /* !_DIDEROT_CL_SCHEDULER_H_ */

Legend:
Removed from v.1501  
changed lines
  Added in v.1513

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