#define PI 3.14159265358979323846
#define PI_2 1.57079632679489661923

__kernel void spinFact(__global float2* w, int n)
{
    unsigned int i = get_global_id(0);

    float2 angle = (float2)(2*i*PI/(float)n,(2*i*PI/(float)n)+PI_2);
    w[i] = cos(angle);
}

__kernel void bitReverse(__global float2 *dst, __global float2 *src, int m, int n)
{
    unsigned int gid = get_global_id(0);
    unsigned int nid = get_global_id(1);

    unsigned int j = gid;
    j = (j & 0x55555555) << 1 | (j & 0xAAAAAAAA) >> 1;
    j = (j & 0x33333333) << 2 | (j & 0xCCCCCCCC) >> 2;
    j = (j & 0x0F0F0F0F) << 4 | (j & 0xF0F0F0F0) >> 4;
    j = (j & 0x00FF00FF) << 8 | (j & 0xFF00FF00) >> 8;
    j = (j & 0x0000FFFF) << 16 | (j & 0xFFFF0000) >> 16;

    j >>= (32-m);

    dst[nid*n+j] = src[nid*n+gid];
}

__kernel void norm(__global float2 *x, int n)
{
    unsigned int gid = get_global_id(0);
    unsigned int nid = get_global_id(1);

    x[nid*n+gid] = x[nid*n+gid] / (float2)((float)n, (float)n);
}

__kernel void butterfly(__global float2 *x, __global float2* w, int m, int n, int iter, uint flag)
{
    unsigned int gid = get_global_id(0);
    unsigned int nid = get_global_id(1);

    int butterflySize = 1 << (iter-1);
    int butterflyGrpDist = 1 << iter;
    int butterflyGrpNum = n >> iter;
    int butterflyGrpBase = (gid >> (iter-1))*(butterflyGrpDist);
    int butterflyGrpOffset = gid & (butterflySize-1);

    int a = nid * n + butterflyGrpBase + butterflyGrpOffset;
    int b = a + butterflySize;

    int l = butterflyGrpNum * butterflyGrpOffset;

    float2 xa, xb, xbxx, xbyy, wab, wayx, wbyx, resa, resb;

    xa = x[a];
    xb = x[b];
    xbxx = xb.xx;
    xbyy = xb.yy;

    wab = as_float2(as_uint2(w[l]) ^ (uint2)(0x0, flag));
    wayx = as_float2(as_uint2(wab.yx) ^ (uint2)(0x80000000, 0x0));
    wbyx = as_float2(as_uint2(wab.yx) ^ (uint2)(0x0, 0x80000000));

    resa = xa + xbxx*wab + xbyy*wayx;
    resb = xa - xbxx*wab + xbyy*wbyx;

    x[a] = resa;
    x[b] = resb;
}

__kernel void transpose(__global float2 *dst, __global float2* src, int n)
{
    unsigned int xgid = get_global_id(0);
    unsigned int ygid = get_global_id(1);

    unsigned int iid = ygid * n + xgid;
    unsigned int oid = xgid * n + ygid;

    dst[oid] = src[iid];
}

__kernel void highPassFilter(__global float2* image, int n, int radius)
{
    unsigned int xgid = get_global_id(0);
    unsigned int ygid = get_global_id(1);

    int2 n_2 = (int2)(n>>1, n>>1);
    int2 mask = (int2)(n-1, n-1);

    int2 gid = ((int2)(xgid, ygid) + n_2) & mask;

    int2 diff = n_2 - gid;
    int2 diff2 = diff * diff;
    int dist2 = diff2.x + diff2.y;

    int2 window;

    if (dist2 < radius*radius) {
        window = (int2)(0L, 0L);
    } else {
        window = (int2)(-1L, -1L);
    }
        image[ygid*n+xgid] = as_float2(as_int2(image[ygid*n+xgid]) & window);
}