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

SCM Repository

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

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

Parent Directory Parent Directory | Revision Log Revision Log


Revision 1489 - (view) (download) (as text)

1 : lamonts 1488 /*! \file cl-scheduler.h
2 :     *
3 :     * \author Lamont Samuels
4 :     */
5 :    
6 :     /*
7 :     * COPYRIGHT (c) 2011 The Diderot Project (http://diderot-language.cs.uchicago.edu)
8 :     * All rights reserved.
9 :     */
10 :    
11 :     #ifndef _DIDEROT_CL_SCHEDULER_H_
12 :     #define _DIDEROT_CL_SCHEDULER_H_
13 :    
14 :     #ifndef _DIDEROT_CL_TYPES_H_
15 :     # include "cl-types.h"
16 :     #endif
17 :    
18 :     /* Scheduler meta-information */
19 :     typedef struct {
20 :     int blkIdx; // the id of this block
21 :     int nActive; // number of active (status != DIE or STABLE) strands
22 :     int nDead; // number of strands in the DIE state
23 :     int nStabilizing; // number of new strands in the STABILIZE state
24 :     int nDying; // number of new strands in the DIE state
25 :     } StrandBlock_t;
26 :    
27 :     typedef struct {
28 :     int numStrands; // number of strands
29 : lamonts 1489 int sId; // the index accumlator for the todo list or queue
30 : lamonts 1488 int nextStrand[1]; // index of the next strand to retrieve from the pool
31 :     int queueSize; // number of blocks on the scheduler's queue
32 :     int todoSize; // number of blocks on the scheduler's todo list
33 :     int numAvailable[1]; // number of active strands left to process
34 :     } SchedState_t;
35 :    
36 :    
37 :     //! \brief Get a block of strands from the scheduler pool.
38 :     //! \param sched global scheduler state
39 :     //! \param blocks global array of strand blocks
40 :     //! \param lBlk pointer to global storage for the strand block
41 :     //! \return lBlk if there is a block to be scheduled, and 0 otherwise
42 :     inline void GetBlock (
43 :     __global SchedState_t *sched,
44 :     __global int * blockIndxs,
45 :     __global int * queue,
46 :     __global StrandBlock_t *blocks,
47 :     __local StrandBlock_t *lBlk)
48 :     {
49 : lamonts 1489 //Retrieve a block from the list
50 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
51 : lamonts 1489
52 :     //Check to see if the index retrieve is less than the
53 :     //number of blocks on the queue
54 : lamonts 1488 if(blkIdx < sched->queueSize)
55 :     {
56 : lamonts 1489 //Copy over the global data for the block into local memory
57 : lamonts 1488 lBlk->blkIdx = (blocks[queue[blkIdx]]).blkIdx;
58 :     lBlk->nActive = (blocks[queue[blkIdx]]).nActive;
59 :     lBlk->nDead = (blocks[queue[blkIdx]]).nDead;
60 :     lBlk->nStabilizing = (blocks[queue[blkIdx]]).nStabilizing;
61 :     lBlk->nDying = (blocks[queue[blkIdx]]).nDying;
62 :    
63 : lamonts 1489 //If the his block has no active strands then check to see if
64 :     // there are strands still waiting to be processed
65 : lamonts 1488 if(lBlk->nActive == 0)
66 :     {
67 : lamonts 1489 //Get a block of strands
68 : lamonts 1488 int startStrandIdx = atom_add(&(sched->nextStrand[0]), BLK_SZ);
69 :     int idx = 0, count = 0;
70 : lamonts 1489 //For each work-item in the workgroup
71 : lamonts 1488 for(int i = startStrandIdx, count = 0; i < sched->numStrands && count < BLK_SZ; i++,count++)
72 :     {
73 : lamonts 1489 //Increment the number of active strands and assign each work-item a strand to process
74 : lamonts 1488 lBlk->nActive++;
75 :     blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;
76 :     idx++;
77 :     }
78 : lamonts 1489 //If there are no strands left to process than make this block inactive
79 : lamonts 1488 if(idx == 0)
80 :     lBlk->nActive = -1;
81 :     }
82 :     }
83 : lamonts 1489 //IF there isn't any queue items then notify the the workgroup
84 : lamonts 1488 else
85 :     lBlk->nActive = -1;
86 :     }
87 :    
88 : lamonts 1489 //! \brief Get a block of strands from the Todolist scheduler pool.
89 : lamonts 1488 //! \param sched global scheduler state
90 :     //! \param blocks global array of strand blocks
91 :     //! \param lBlk pointer to global storage for the strand block
92 :     //! \return lBlk if there is a block to be scheduled, and 0 otherwise
93 :     inline void GetBlockFromTodoList (
94 :     __global SchedState_t *sched,
95 :     __global int * blockIndxs,
96 :     __global int * todoList,
97 :     __global StrandBlock_t *blocks,
98 :     __local StrandBlock_t *lBlk)
99 :     {
100 : lamonts 1489 //Retrieve a block from the todolist
101 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
102 :    
103 : lamonts 1489 //Check to make sure the idx retrieve is a valid for the list
104 : lamonts 1488 if(blkIdx < sched->todoSize)
105 :     {
106 : lamonts 1489 //Copy over the global data into the local data.
107 : lamonts 1488 lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;
108 :     lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;
109 :     lBlk->nDead = (blocks[todoList[blkIdx]]).nDead;
110 :     lBlk->nStabilizing = (blocks[todoList[blkIdx]]).nStabilizing;
111 :     lBlk->nDying = (blocks[todoList[blkIdx]]).nDying;
112 :     }
113 :     else
114 :     lBlk->nActive = -1;
115 :     }
116 :    
117 :     // Exclusive Scan that requires the number of work-items to be less than
118 :     // the maximum number of work-items for a single work-group And the only
119 :     // one work group to be executed.
120 :     inline void scan ( __global int * input,
121 :     __global int * status,
122 :     __local int * output,
123 :     StrandStatus_t statusTy,
124 :     int n,
125 :     __local int * temp)
126 :     {
127 :     int thid = get_local_id(0);
128 :     int pout = 0, pin = 1;
129 :    
130 :     //load input into local memory
131 :     // Exclusive scan: shift right by one and set first element to 0
132 :     if(thid == 0 || status[(*(input + thid))] != statusTy)
133 :     temp[thid] = 0;
134 :     else
135 :     temp[thid] = 1;
136 :    
137 :     barrier(CLK_LOCAL_MEM_FENCE);
138 :    
139 :     for(int offset = 1; offset < n; offset *= 2)
140 :     {
141 :     pout = 1 - pout;
142 :     pin = 1 - pout;
143 :    
144 :     if(thid >= offset)
145 :     temp[pout * n + thid] = temp[pin * n + thid] + temp[pin * n + thid - offset];
146 :     else
147 :     temp[pout * n + thid] = temp[pin * n + thid];
148 :    
149 :     barrier(CLK_LOCAL_MEM_FENCE);
150 :     }
151 :     output[thid] = temp[pout* n + thid];
152 :     }
153 :     #endif /* !_DIDEROT_CL_SCHEDULER_H_ */

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