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

SCM Repository

[diderot] View of /trunk/src/compiler/cl-target/fragments/update.in
ViewVC logotype

View of /trunk/src/compiler/cl-target/fragments/update.in

Parent Directory Parent Directory | Revision Log Revision Log


Revision 2000 - (download) (annotate)
Thu Oct 4 18:07:32 2012 UTC (6 years, 10 months ago) by jhr
File size: 2214 byte(s)
  fix spelling error
__kernel __attribute__((reqd_work_group_size(BLK_SZ, 1, 1))) 
void Diderot_UpdateKernel (
        __global @STRAND_TY@ *stateIn, 
        __global @STRAND_TY@ *stateOut, 
        __global int *status, 
        __global SchedState_t *sched, 
        __global StrandBlock_t *blocks, 
        __global int *blockIndxs, 
        __global int *queue, 
        __global int *todoList, 
        __local StrandBlock_t *bp@SHOULD_ADD_GLOBALS@
        @GLOBALS@@SHOULD_ADD_DATA@
        @GLOBAL_DATA_PTRS@)
{
    @GLOBAL_DATA_ASSIGN@ 

    int id = get_local_id(0);
    do {
        if (id == 0) {
            GetBlock (sched, blockIndxs, queue, blocks, bp);
            if (bp->nActive!=-1) {
                bp->nStabilizing = 0;
                bp->nDying = 0;
            }
        }
        barrier(CLK_LOCAL_MEM_FENCE);
        if (bp->nActive != -1) {
            int idx = blockIndxs[bp->blkIdx*BLK_SZ+id];
            if (id < bp->nActive) {
                int sts = Diderot_Update(&stateIn[idx], &stateOut[idx], @GLOBAL_VAR@, @DATA_VAR@);
                switch (sts) {
                  case DIDEROT_DIE:
                    atom_inc (&bp->nDying);
                    status[idx] = DIDEROT_DIE;
                    atom_dec (&sched->numAvailable);
                    break;
                  case DIDEROT_STABILIZE:
                    atom_inc (&bp->nStabilizing);
                    status[idx] = DIDEROT_STABLE;
                    atom_dec (&sched->numAvailable);
		    Diderot_Stabilize(&stateOut[idx], &stateOut[idx], @GLOBAL_VAR@, @DATA_VAR@);
                    break;
                  case DIDEROT_ACTIVE:
                    break;
                }
                //Copy out-state into the in-state for the next iteration 
                Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
            }
            barrier(CLK_LOCAL_MEM_FENCE);
            if (id == 0) {
                //Copy from local to global memory 
                StrandBlock_Copy(&blocks[bp->blkIdx],bp); 
                //Add Block to the todo list 
                int todoIdx = atom_inc(&sched->todoSize);
                todoList[todoIdx] = bp->blkIdx;
            }
        }
    } while (bp->nActive != -1);
}

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