Fill a 2D array with pyopenCL

Question

I am trying to fill a 2D array using pyOpenCL. The compute kernel and its call are posted below:

ctx = cl.Context([cl.get_platforms()[0].get_devices()[0]])
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
x_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=x)
y_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=y)
a_buf = cl.Buffer(ctx, mf.WRITE_ONLY, a.nbytes)

prg = cl.Program(ctx, """
//#define PYOPENCL_DEFINE_CDOUBLE
#include "pyopencl-complex.h"    
__kernel void makeA(const unsigned int ySize, const float cov,
                  const int x0, const int y0, __global const float *x, __global const float *y,
                  __global cfloat_t *a)
{
  int gid0 = get_global_id(0);
  int gid1 = get_global_id(1);

  a[gid1 + ySize*gid0] = (cfloat_t)(1, 0);
}
""").build()

prg.makeA(queue, a.shape, None, np.int32(ySize),
    np.float32(c), np.int32(x0), np.int32(y0), x_buf, y_buf, a_dest_buf)  
cl.enqueue_copy(queue, a, a_dest_buf)

Now, this seems to work alright. Today I learned, that it is very usefull to use workgroups. This, I do not get to work properly. I have tried to replace the function call by

prg.makeA(queue, a.shape, (16,16), np.int32(ySize),
    np.float32(c), np.int32(x0), np.int32(y0), x_buf, y_buf, a_dest_buf)

but I do not know how to properly compute the new x and y indicies of the array and, therefore, I cannot change

a[yIdx + ySize*xIdx] = (cfloat)( x[xIdx] , 0);  

Answer 1

I now tried the following, which worked:

__kernel void makea(const unsigned int ySize, const float cov,
                  const int x0, const int y0, __global const float *x, __global const float *y,
                  __global cfloat_t *a)
{

  int xIdx = get_local_id(0)+get_group_id(0)*get_local_size(0);
  int yIdx = get_local_id(1)+get_group_id(1)*get_local_size(1);

  a[yIdx + ySize*xIdx] = (cfloat_t)(1, 0);
}
""").build()




prg.makeA(queue, a.shape, (32,32), np.int32(ySize),
    np.float32(c), np.int32(x0), np.int32(y0), x_buf, y_buf, a_dest_buf)

By this I introduced a workgroup of size 32x32 [function call parameter: (32,32)]. It appears that this now makes use of several multi stream processors.