code: 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
docs: 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.