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

SCM Repository

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

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

Parent Directory Parent Directory | Revision Log Revision Log


Revision 1492 - (download) (as text) (annotate)
Fri Sep 9 13:13:47 2011 UTC (8 years ago) by jhr
File size: 5613 byte(s)
  style fixes
/*! \file cl-scheduler.h
 *
 * \author Lamont Samuels
 */

/*
 * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)
 * All rights reserved.
 */

#ifndef _DIDEROT_CL_SCHEDULER_H_
#define _DIDEROT_CL_SCHEDULER_H_
 
#ifndef _DIDEROT_CL_TYPES_H_
#  include "cl-types.h"
#endif

/* 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[1];              // index of the next strand to retrieve from the pool
    int queueSize;                  // number of blocks on the scheduler's queue  
    int todoSize;                   // number of blocks on the scheduler's todo list
    int numAvailable[1];            // 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[0]), 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 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))] != 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]; 
}
#endif /* !_DIDEROT_CL_SCHEDULER_H_ */

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