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 1989 - (download) (annotate)
Fri Aug 3 01:54:21 2012 UTC (8 years, 9 months ago) by lamonts
File size: 2227 byte(s)
Removed old code that caused problems on AMD GPUs and fixed the type conversion problems that caused problems for Mandelbrot. the GPU backend calls stabalize properly now. 
__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@

    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;
        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);
                  case DIDEROT_STABILIZE:
                    atom_inc (&bp->nStabilizing);
                    status[idx] = DIDEROT_STABLE;
                    atom_dec (&sched->numAvailable);
                    Diderot_Stabailize(&stateOut[idx],&stateOut[idx],@GLOBAL_VAR@, @DATA_VAR@);
                  case DIDEROT_ACTIVE:
                //Copy out-state into the in-state for the next iteration 
                Diderot_StateCopy (&stateOut[idx], &stateIn[idx]);
            if (id == 0) {
                //Copy from local to global memory 
                //Add Block to the todo list 
                int todoIdx = atom_inc(&sched->todoSize);
                todoList[todoIdx] = bp->blkIdx;
    } while (bp->nActive != -1);

ViewVC Help
Powered by ViewVC 1.0.0