Проблема с индексацией PyOpenCL

Я экспериментирую с OpenCl в python и не могу понять, что я делаю неправильно с этим простым кодом копирования матрицы.

Моя входная матрица:
[[1 2 3 4],
[5 6 7 8],
[9 10 11 12],
[13 14 15 16],
[ 17 18 19 20]]

Я получаю следующий вывод:
[[1 2 3 4],
[5 6 7 8],
[9 10 0 0],
[0 0 0 0],
[0 0 0 0]]

Почему копируется только часть моей матрицы? Что я делаю не так?

Вот мой код:

import pyopencl as cl
import numpy as np

kernel = """
__kernel void
copy( __global const float *g_data, const int h, const int w, __global float *g_out )
{
// Get global position
size_t row = get_global_id(0);
const int s = row * w;
__global const float *in = &g_data[ s ];
__global float *out = &g_out[ s ];
for(int i=0; i<w; ++i)
{
          out[i] = in[i];
}
}
"""

class test:
      def __init__(self):
                 # Create opencl context
                 platform = cl.get_platforms()[2]
                 self.__ctx__ = cl.Context( [platform.get_devices()[0]] )
                 # Create opencl queue
                 self.__queue__ = cl.CommandQueue(self.__ctx__)
                 # Build opencl kernel
                 self.__kernel__ = cl.Program(self.__ctx__, kernel).build()


      def __del__(self):
             del self.__queue__
             del self.__kernel__
             del self.__ctx__

      def __call__(self, data):
                 # Get matrix dimensions
                 h, w = data.shape
                 mf = cl.mem_flags
                 # Set input buffer
                 g_data = cl.Buffer(self.__ctx__, (mf.READ_ONLY | mf.COPY_HOST_PTR), hostbuf=data)
                 # Set output buffer
                 self.__out__ = np.zeros( data.shape, dtype=np.float )
                 g_out = cl.Buffer(self.__ctx__, mf.WRITE_ONLY, self.__out__.nbytes)
                 # Run kernel
                 kernel_event = self.__kernel__.copy(
                                                                    self.__queue__,
                                                                    (h,),
                                                                    None,
                                                                    g_data,
                                                                    np.int32(h),
                                                                    np.int32(w),
                                                                    g_out,
                                                                    wait_for=None
                                                                )
                 # Copy data
                 out_event = cl.enqueue_copy(self.__queue__, self.__out__, g_out, wait_for=[kernel_event])
                 out_event.wait()
                 # Free memory
                 g_out.release()
                 print( self.__out__ )