extern "C" {
    __global__ void
    swap(long* param, int* Data)
    {
        long block = param[0];
        long first = param[1];
        int x = blockIdx.x*blockDim.x+threadIdx.x;
        
        int position = x/block;
        int index = x+block*position;
        
        block = (first == 1) ? ((block<<1)*(position+1))-(index%block)-1 : index+block;
    
        if (Data[block] < Data[index]) {
            int tmp = Data[index];
            Data[index] = Data[block];
            Data[block] = tmp;
        }
    }        
}