How to pass an array of vectors in pyOpenCL
===

I'm moving a simulation into pyOpenCL and can't get my data access to work. I'm trying to supply a 1D array of vectors (well, actually several, but the example I've included just used one).

Currently, several vectors are copied over just fine, but then the data is simply not what I supplied.

I don't think I've posted here before, so apologies if any of the formatting/presentation is wrong. Also, I've just stripped out all the simulation code, so I realise this code is currently not actually doing anything, I just want to get the buffer passing correct.

Thanks in advance.

The kernel (kertest.py):

    step1 = """
    #pragma OPENCL EXTENSION cl_amd_printf: enable
    #define X xdim
    #define Y ydim
    __kernel void k1(__global float3 *spins,
                     __local float3 *tile)
    {        
        ushort lid = 2 * get_local_id(0);
        ushort group = 2 * get_group_id(0);
        ushort num = get_num_groups(0);
        int lim = X*Y*3;
        
        for (ushort i = 0; i < lim; i++)
            {
                if (lid == 0 && group == 0)
                {
                    printf("%f :: %d\\n", spins[i].x, i);
                }
             }
    }"""

The code itself (gputest.py):

    import kertest as k2D
    import numpy as np
    import pyopencl as cl
    
    class GPU_MC2DSim():
        def __init__(self, x, y):
            self.x = x
            self.y = y
            
            if x >= y:
                self.xdim = int(self.x)
                self.ydim = int(self.y)
            else:
                self.xdim = int(self.y)
                self.ydim = int(self.x)
            
            if self.xdim % 2 != 0: self.xdim += 1
            
            if self.ydim % 2 != 0: self.ydim += 1
                    
            self.M = np.ones((self.xdim*self.ydim, 3)).astype(np.float32)
            self.M[:, 1] += 1.0
            self.M[:, 2] += 2.0
            
            print self.M
            
        def simulate(self):
            ctx = cl.create_some_context()
            q = cl.CommandQueue(ctx)
            mf = cl.mem_flags
            
            #Pass buffer:
            M_buf = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf = self.M)
            
            #Insert kernel parameters:
            params = {"xdim" : "%d" % (self.xdim),
                      "ydim" : "%d" % (self.ydim),
                      }
            for name in params:
                k2D.step1 = k2D.step1.replace(name, params[name])
                
            #Compile kernel:
            step1 = cl.Program(ctx, k2D.step1).build()
            
            locmem = cl.LocalMemory(self.xdim*4*4)
            
            step1.k1(q, ((self.xdim*self.ydim)/4,), (self.xdim/2,), M_buf, locmem).wait()
            return None
        
    xdim = 4
    ydim = 4
    sim = GPU_MC2DSim(xdim, ydim)
    sim.simulate()


