extern "C" {
    
    __global__ void
    bitReverse(long* param, float* src, float* dst)
    {
        unsigned long gid = blockIdx.x*blockDim.x+threadIdx.x; // (unsigned long)s->get_param(0);
        unsigned long nid = blockIdx.y*blockDim.y+threadIdx.y; // (unsigned long)s->get_param(1);
        
        unsigned int j = gid;
        
        unsigned long m = param[0];
        unsigned long n = param[1];
        
        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)*2] = src[(nid*n+gid)*2];
        dst[(nid*n+j)*2+1] = src[(nid*n+gid)*2+1];
    }
}