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

SCM Repository

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

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

Parent Directory Parent Directory | Revision Log Revision Log


Revision 1543 - (download) (annotate)
Mon Oct 17 19:53:03 2011 UTC (10 years ago) by jhr
File size: 1981 byte(s)
  Code cleanup
__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, 
        __global Diderot_Globals_t *diderotGlobals, 
        @GLOBAL_DATA_PTRS@)
{
    Diderot_data_ptr_t diderotDataPtrs;
    @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], diderotGlobals, diderotDataPtrs);
                switch (sts) {
                  case DIDEROT_DIE:
                    atom_inc (&blocks[bp->blkIdx].nDying);
                    status[idx] = sts;
                    break;
                  case DIDEROT_STABILIZE:
                    atom_inc (&blocks[bp->blkIdx].nStabilizing);
                    status[idx] = DIDEROT_STABLE;
                    Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
                    atom_dec (&sched->numAvailable);
                    break;
                  case DIDEROT_ACTIVE:
                    Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
                    break;
                }
            }
            barrier(CLK_LOCAL_MEM_FENCE);
            if (id == 0) {
                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