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

SCM Repository

[diderot] Annotation of /branches/pure-cfg/src/lib/cl-target/cl-kernels/update_cl.in
ViewVC logotype

Annotation of /branches/pure-cfg/src/lib/cl-target/cl-kernels/update_cl.in

Parent Directory Parent Directory | Revision Log Revision Log


Revision 1642 - (view) (download)

1 : lamonts 1541 __kernel __attribute__((reqd_work_group_size(BLK_SZ, 1, 1)))
2 : jhr 1543 void Diderot_UpdateKernel (
3 :     __global @STRAND_TY@ *stateIn,
4 :     __global @STRAND_TY@ *stateOut,
5 :     __global int *status,
6 :     __global SchedState_t *sched,
7 :     __global StrandBlock_t *blocks,
8 :     __global int *blockIndxs,
9 : lamonts 1642 __global int *queue,
10 :     __global int *todoList,
11 : lamonts 1641 __local StrandBlock_t *bp@SHOULD_ADD_GLOBALS@
12 :     @GLOBALS@@SHOULD_ADD_DATA@
13 : jhr 1543 @GLOBAL_DATA_PTRS@)
14 : lamonts 1541 {
15 :     @GLOBAL_DATA_ASSIGN@
16 :    
17 :     int id = get_local_id(0);
18 : jhr 1543 do {
19 :     if (id == 0) {
20 :     GetBlock (sched, blockIndxs, queue, blocks, bp);
21 : lamonts 1541 if (bp->nActive!=-1) {
22 :     bp->nStabilizing = 0;
23 :     bp->nDying = 0;
24 :     }
25 : jhr 1543 }
26 : lamonts 1541 barrier(CLK_LOCAL_MEM_FENCE);
27 : jhr 1543 if (bp->nActive != -1) {
28 : lamonts 1541 int idx = blockIndxs[bp->blkIdx*BLK_SZ+id];
29 : jhr 1543 if (id < bp->nActive) {
30 : lamonts 1642 int sts = Diderot_Update(&stateIn[idx], &stateOut[idx], @GLOBAL_VAR@, @DATA_VAR@);
31 : jhr 1543 switch (sts) {
32 :     case DIDEROT_DIE:
33 :     atom_inc (&blocks[bp->blkIdx].nDying);
34 : lamonts 1541 status[idx] = sts;
35 : jhr 1543 break;
36 :     case DIDEROT_STABILIZE:
37 :     atom_inc (&blocks[bp->blkIdx].nStabilizing);
38 : lamonts 1541 status[idx] = DIDEROT_STABLE;
39 : jhr 1543 Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
40 :     atom_dec (&sched->numAvailable);
41 :     break;
42 :     case DIDEROT_ACTIVE:
43 :     Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
44 :     break;
45 :     }
46 :     }
47 : lamonts 1541 barrier(CLK_LOCAL_MEM_FENCE);
48 : jhr 1543 if (id == 0) {
49 : lamonts 1541 int todoIdx = atom_inc(&sched->todoSize);
50 :     todoList[todoIdx] = bp->blkIdx;
51 :     }
52 :     }
53 : jhr 1543 } while (bp->nActive != -1);
54 : lamonts 1541 }

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