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 1501 - (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 : jhr 1492 int blkIdx; // the id of this block
21 : lamonts 1488 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 : jhr 1492 int sId; // the index accumlator for the todo list or queue
30 : lamonts 1500 int nextStrand; // index of the next strand to retrieve from the pool
31 : lamonts 1501 int clearQueueSz; // an indicator on whether the queue size should be cleared
32 : lamonts 1488 int queueSize; // number of blocks on the scheduler's queue
33 :     int todoSize; // number of blocks on the scheduler's todo list
34 : lamonts 1500 int numAvailable; // number of active strands left to process
35 : lamonts 1488 } SchedState_t;
36 :    
37 :    
38 :     //! \brief Get a block of strands from the scheduler pool.
39 :     //! \param sched global scheduler state
40 :     //! \param blocks global array of strand blocks
41 :     //! \param lBlk pointer to global storage for the strand block
42 :     //! \return lBlk if there is a block to be scheduled, and 0 otherwise
43 :     inline void GetBlock (
44 : jhr 1492 __global SchedState_t *sched,
45 :     __global int *blockIndxs,
46 :     __global int *queue,
47 :     __global StrandBlock_t *blocks,
48 :     __local StrandBlock_t *lBlk)
49 : lamonts 1488 {
50 : jhr 1492 // Retrieve a block from the list
51 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
52 : lamonts 1489
53 : jhr 1492 // Check to see if the index retrieve is less than the
54 :     // number of blocks on the queue
55 :     if (blkIdx < sched->queueSize) {
56 :     // 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 : jhr 1492 // If the his block has no active strands then check to see if
64 :     // there are strands still waiting to be processed
65 :     if (lBlk->nActive == 0) {
66 :     // Get a block of strands
67 : lamonts 1500 int startStrandIdx = atom_add(&sched->nextStrand, BLK_SZ);
68 : lamonts 1488 int idx = 0, count = 0;
69 : jhr 1492 // For each work-item in the workgroup
70 :     for (int i = startStrandIdx, count = 0; i < sched->numStrands && count < BLK_SZ; i++,count++) {
71 :     // Increment the number of active strands and assign each work-item a strand to process
72 : lamonts 1488 lBlk->nActive++;
73 :     blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;
74 :     idx++;
75 :     }
76 : jhr 1492 // If there are no strands left to process than make this block inactive
77 : lamonts 1488 if(idx == 0)
78 :     lBlk->nActive = -1;
79 :     }
80 : jhr 1492 }
81 :     else // if there isn't any queue items then notify the the workgroup
82 : lamonts 1488 lBlk->nActive = -1;
83 :     }
84 : lamonts 1501 //! \brief Copies the strand black from __local space to __global space
85 :     //! \param sched global scheduler state
86 :     //! \param blocks global array of strand blocks
87 :     //! \param lBlk pointer to global storage for the strand block
88 :     //! \return lBlk if there is a block to be scheduled, and 0 otherwise
89 :     inline void StrandBlock_Copy (__global StrandBlock_t * gBlock,
90 :     __local StrandBlock_t * lBlock)
91 :     {
92 :     gBlock->nActive = lBlock->nActive;
93 :     gBlock->nDead = lBlock->nDead;
94 :     gBlock->nStabilizing = lBlock->nStabilizing;
95 :     gBlock->nDying = lBlock->nDying;
96 :     }
97 : lamonts 1489 //! \brief Get a block of strands from the Todolist scheduler pool.
98 : lamonts 1488 //! \param sched global scheduler state
99 :     //! \param blocks global array of strand blocks
100 :     //! \param lBlk pointer to global storage for the strand block
101 :     //! \return lBlk if there is a block to be scheduled, and 0 otherwise
102 :     inline void GetBlockFromTodoList (
103 :     __global SchedState_t *sched,
104 : jhr 1492 __global int *blockIndxs,
105 :     __global int *todoList,
106 : lamonts 1488 __global StrandBlock_t *blocks,
107 :     __local StrandBlock_t *lBlk)
108 :     {
109 : jhr 1492 // Retrieve a block from the todolist
110 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
111 :    
112 : jhr 1492 // Check to make sure the idx retrieve is a valid for the list
113 :     if (blkIdx < sched->todoSize) {
114 :     // Copy over the global data into the local data.
115 : lamonts 1488 lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;
116 :     lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;
117 :     lBlk->nDead = (blocks[todoList[blkIdx]]).nDead;
118 :     lBlk->nStabilizing = (blocks[todoList[blkIdx]]).nStabilizing;
119 : lamonts 1501 lBlk->nDying = (blocks[todoList[blkIdx]]).nDying;
120 : lamonts 1488 }
121 :     else
122 :     lBlk->nActive = -1;
123 :     }
124 :    
125 : jhr 1492 // Exclusive Scan that requires the number of work-items to be less than
126 :     // the maximum number of work-items for a single work-group and the only
127 :     // one work group to be executed.
128 :     inline void scan (
129 :     __global int *input,
130 :     __global int *status,
131 :     __local int *output,
132 :     StrandStatus_t statusTy,
133 :     int n,
134 :     __local int *temp)
135 :     {
136 :     int thid = get_local_id(0);
137 :     int pout = 0, pin = 1;
138 :    
139 :     // load input into local memory
140 :     // Exclusive scan: shift right by one and set first element to 0
141 : lamonts 1500 if (thid == 0 || status[(*(input + thid - 1))] != statusTy)
142 : jhr 1492 temp[thid] = 0;
143 :     else
144 :     temp[thid] = 1;
145 :    
146 :     barrier(CLK_LOCAL_MEM_FENCE);
147 :    
148 :     for (int offset = 1; offset < n; offset *= 2) {
149 :     pout = 1 - pout;
150 :     pin = 1 - pout;
151 :    
152 :     if(thid >= offset)
153 :     temp[pout * n + thid] = temp[pin * n + thid] + temp[pin * n + thid - offset];
154 :     else
155 :     temp[pout * n + thid] = temp[pin * n + thid];
156 :    
157 :     barrier(CLK_LOCAL_MEM_FENCE);
158 :     }
159 :     output[thid] = temp[pout* n + thid];
160 : lamonts 1488 }
161 :     #endif /* !_DIDEROT_CL_SCHEDULER_H_ */

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