__kernel
void bitReverse(__constant long *param, __global float2 *src,__global float2 *dst)
{
    unsigned long gid = (unsigned long)get_global_id(0);
    unsigned long nid = (unsigned long)get_global_id(1);
    
    unsigned long m = (unsigned long)param[0];
    unsigned long n = (unsigned long)param[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];
}