def contiguous_slice(in1):
    """
    This function unpads an array on the GPU in such a way as to make it contiguous.

    INPUTS:
    in1     (no default):   Array containing data which has been padded.

    OUTPUTS:
    gpu_out1                Array containing unpadded, contiguous data.
    """

    ker = SourceModule("""
                        __global__ void contiguous_slice_ker(float *in1, float *out1)
                        {
                            const int len = gridDim.x*blockDim.x;
                            const int col = (blockDim.x * blockIdx.x + threadIdx.x);
                            const int row = (blockDim.y * blockIdx.y + threadIdx.y);
                            const int tid2 = col + len*row;

                            const int first_idx = len/4;
                            const int last_idx = (3*len)/4;

                            const int out_idx = (col-first_idx)+(row-first_idx)*(len/2);

                            if (((col>=first_idx)&(row>=first_idx))&((col<last_idx)&(row<last_idx)))
                                { out1[out_idx] = in1[tid2]; }

                        }
                       """, keep=True)

    gpu_out1 = gpuarray.empty([in1.shape[0]/2,in1.shape[1]/2], np.float32)

    contiguous_slice_ker = ker.get_function("contiguous_slice_ker")
    contiguous_slice_ker(in1, gpu_out1, block=(32,32,1), grid=(int(in1.shape[1]//32), int(in1.shape[0]//32)))

    return gpu_out1
def contiguous_slice ( in1 ) : ker = SourceModule ( """
                        __global__ void contiguous_slice_ker(float *in1, float *out1)
                        {
                            const int len = gridDim.x*blockDim.x;
                            const int col = (blockDim.x * blockIdx.x + threadIdx.x);
                            const int row = (blockDim.y * blockIdx.y + threadIdx.y);
                            const int tid2 = col + len*row;
                            const int first_idx = len/4;
                            const int last_idx = (3*len)/4;
                            const int out_idx = (col-first_idx)+(row-first_idx)*(len/2);
                            if (((col>=first_idx)&(row>=first_idx))&((col<last_idx)&(row<last_idx)))
                                { out1[out_idx] = in1[tid2]; }
                        }
                       """ , keep = True ) gpu_out1 = gpuarray . empty ( [ in1 . shape [ 0 ] / 2 , in1 . shape [ 1 ] / 2 ] , np . float32 ) contiguous_slice_ker = ker . get_function ( "contiguous_slice_ker" ) contiguous_slice_ker ( in1 , gpu_out1 , block = ( 32 , 32 , 1 ) , grid = ( int ( in1 . shape [ 1 ] // 32 ) , int ( in1 . shape [ 0 ] // 32 ) ) ) return gpu_out1
This function unpads an array on the GPU in such a way as to make it contiguous.

    INPUTS:
    in1     (no default):   Array containing data which has been padded.

    OUTPUTS:
    gpu_out1                Array containing unpadded, contiguous data.