permute_tag.inl

/* -----------------------------------------------------------------------------
 *
 * Module    : Permute
 * Copyright : (c) [2009..2010] Trevor L. McDonell
 * License   : BSD
 *
 * Forward permutation, characterised by a function that determines for each
 * element in the source array where it should go in the target. The output
 * array should be initialised with a default value, as the permutation may be
 * between arrays of different sizes and some positions may never be touched.
 *
 * Elements from the source array are dropped for which the permutation function
 * yields the magic index `ignore`.
 *
 * ---------------------------------------------------------------------------*/


#define TAG_MASK        ((1 << 27) - 1)
#define TAG_THREAD      (threadIdx.x << 27)

extern "C"
__global__ void
permute
(
    ArrOut              d_out,
    const ArrIn0        d_in0,
    const Ix            shape
)
{
    Ix       dst;
    Ix       idx;
    const Ix gridSize = __umul24(blockDim.x, gridDim.x);

    for (idx = __umul24(blockDim.x, blockIdx.x) + threadIdx.x; idx < shape; idx += gridSize)
    {
        dst = project(idx);

        if (dst != ignore)
        {
            TyOut x1;
            TyIn0 x0 = get0(d_in0, idx);

            do
            {
                x1 = get0(d_out, dst) & TAG_MASK;
                x1 = apply(x0, x1)    | TAG_THREAD;
                set(d_out, dst, x1);

                __syncthreads();
            }
            while (get0(d_out, dst) != x1);

            __syncthreads();
            set(d_out, dst, get0(d_out, dst) & TAG_MASK);
        }
    }
}