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 1500 - (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 1488 int queueSize; // number of blocks on the scheduler's queue
32 :     int todoSize; // number of blocks on the scheduler's todo list
33 : lamonts 1500 int numAvailable; // number of active strands left to process
34 : lamonts 1488 } 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 : jhr 1492 __global SchedState_t *sched,
44 :     __global int *blockIndxs,
45 :     __global int *queue,
46 :     __global StrandBlock_t *blocks,
47 :     __local StrandBlock_t *lBlk)
48 : lamonts 1488 {
49 : jhr 1492 // Retrieve a block from the list
50 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
51 : lamonts 1489
52 : jhr 1492 // Check to see if the index retrieve is less than the
53 :     // number of blocks on the queue
54 :     if (blkIdx < sched->queueSize) {
55 :     // Copy over the global data for the block into local memory
56 : lamonts 1488 lBlk->blkIdx = (blocks[queue[blkIdx]]).blkIdx;
57 :     lBlk->nActive = (blocks[queue[blkIdx]]).nActive;
58 :     lBlk->nDead = (blocks[queue[blkIdx]]).nDead;
59 :     lBlk->nStabilizing = (blocks[queue[blkIdx]]).nStabilizing;
60 :     lBlk->nDying = (blocks[queue[blkIdx]]).nDying;
61 :    
62 : jhr 1492 // If the his block has no active strands then check to see if
63 :     // there are strands still waiting to be processed
64 :     if (lBlk->nActive == 0) {
65 :     // Get a block of strands
66 : lamonts 1500 int startStrandIdx = atom_add(&sched->nextStrand, BLK_SZ);
67 : lamonts 1488 int idx = 0, count = 0;
68 : jhr 1492 // For each work-item in the workgroup
69 :     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
71 : lamonts 1488 lBlk->nActive++;
72 :     blockIndxs[lBlk->blkIdx * BLK_SZ + idx] = i;
73 :     idx++;
74 :     }
75 : jhr 1492 // If there are no strands left to process than make this block inactive
76 : lamonts 1488 if(idx == 0)
77 :     lBlk->nActive = -1;
78 :     }
79 : jhr 1492 }
80 :     else // if there isn't any queue items then notify the the workgroup
81 : lamonts 1488 lBlk->nActive = -1;
82 :     }
83 :    
84 : lamonts 1489 //! \brief Get a block of strands from the Todolist scheduler pool.
85 : lamonts 1488 //! \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 GetBlockFromTodoList (
90 :     __global SchedState_t *sched,
91 : jhr 1492 __global int *blockIndxs,
92 :     __global int *todoList,
93 : lamonts 1488 __global StrandBlock_t *blocks,
94 :     __local StrandBlock_t *lBlk)
95 :     {
96 : jhr 1492 // Retrieve a block from the todolist
97 : lamonts 1488 int blkIdx = atom_inc(&(sched->sId));
98 :    
99 : jhr 1492 // Check to make sure the idx retrieve is a valid for the list
100 :     if (blkIdx < sched->todoSize) {
101 :     // Copy over the global data into the local data.
102 : lamonts 1488 lBlk->blkIdx = (blocks[todoList[blkIdx]]).blkIdx;
103 :     lBlk->nActive = (blocks[todoList[blkIdx]]).nActive;
104 :     lBlk->nDead = (blocks[todoList[blkIdx]]).nDead;
105 :     lBlk->nStabilizing = (blocks[todoList[blkIdx]]).nStabilizing;
106 :     lBlk->nDying = (blocks[todoList[blkIdx]]).nDying;
107 :     }
108 :     else
109 :     lBlk->nActive = -1;
110 :     }
111 :    
112 : jhr 1492 // Exclusive Scan that requires the number of work-items to be less than
113 :     // the maximum number of work-items for a single work-group and the only
114 :     // one work group to be executed.
115 :     inline void scan (
116 :     __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 : lamonts 1500 if (thid == 0 || status[(*(input + thid - 1))] != statusTy)
129 : jhr 1492 temp[thid] = 0;
130 :     else
131 :     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 : lamonts 1488 }
148 :     #endif /* !_DIDEROT_CL_SCHEDULER_H_ */

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