title: Struct Alignment with PyOpenCL
body: <p><strong>update: the <code>int4</code> in my kernel was wrong.</strong></p>
  
  <p>I am using pyopencl but am unable to get struct alignment to work correctly.  In the code below, which calls the kernel twice, the <code>b</code> value is returned correctly (as 1), but the <code>c</code> value has some "random" value.</p>
  
  <p>In other words: <strong><em>I am trying to read two members of a struct.  I can read the first but not the second.  Why?</em></strong></p>
  
  <p>The same issue occurs whether I use numpy structured arrays or pack with struct.  And the <code>_-attribute__</code> settings in the comments don't help either.</p>
  
  <p>I suspect I am doing something stupid elsewhere in the code, but can't see it.  Any help appreciated.</p>
  
  <pre><code>import struct as s
  import pyopencl as cl
  import numpy as n
  
  ctx = cl.create_some_context()
  queue = cl.CommandQueue(ctx)
  
  for use_struct in (True, False):
  
      if use_struct:
          a = s.pack('=ii',1,2)
          print(a, len(a))
          a_dev = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, len(a))
      else:
  #       a = n.array([(1,2)], dtype=n.dtype('2i4', align=True))
          a = n.array([(1,2)], dtype=n.dtype('2i4'))
          print(a, a.itemsize, a.nbytes)
          a_dev = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, a.nbytes)
  
      b = n.array([0], dtype='i4')
      print(b, b.itemsize, b.nbytes)
      b_dev = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, b.nbytes)
  
      c = n.array([0], dtype='i4')
      print(c, c.itemsize, c.nbytes)
      c_dev = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, c.nbytes)
  
      prg = cl.Program(ctx, """
          typedef struct s {
              int4 f0;
              int4 f1 __attribute__ ((packed));
  //            int4 f1 __attribute__ ((aligned (4)));
  //            int4 f1;
          } s;
          __kernel void test(__global const s *a, __global int4 *b, __global int4 *c) {
              *b = a-&gt;f0;
              *c = a-&gt;f1;
          }
          """).build()
  
      cl.enqueue_copy(queue, a_dev, a)
      event = prg.test(queue, (1,), None, a_dev, b_dev, c_dev)
      event.wait()
      cl.enqueue_copy(queue, b, b_dev)
      print(b)
      cl.enqueue_copy(queue, c, c_dev)
      print(c)
  </code></pre>
  
  <p>The output (I had to reformat while cut+pasting, so may have messed up line breaks slightly; I've also added comments indicating what the various print values are):</p>
  
  <pre><code># first using struct
  /home/andrew/projects/personal/kultrung/env/bin/python3.2 /home/andrew/projects/personal/kultrung/src/kultrung/test6.py
  b'\x01\x00\x00\x00\x02\x00\x00\x00' 8 # the struct packed values
  [0] 4 4                               # output buffer 1
  [0] 4 4                               # output buffer 2
  /home/andrew/projects/personal/kultrung/env/lib/python3.2/site-packages/pyopencl/cache.py:343: UserWarning: Build succeeded, but resulted in non-empty logs: Build on &lt;pyopencl.Device 'Intel(R) Core(TM)2 CPU         T5600  @ 1.83GHz' at 0x1385a20&gt; succeeded, but said:
  
  Build started Kernel &lt;test&gt; was successfully vectorized Done.   warn("Build succeeded, but resulted in non-empty logs:\n"+message)
  [1]         # the first value (correct)
  [240]       # the second value (wrong)
  
  # next using numpy
  [[1 2]] 4 8 # the numpy struct
  [0] 4 4     # output buffer
  [0] 4 4     # output buffer
  /home/andrew/projects/personal/kultrung/env/lib/python3.2/site-packages/pyopencl/__init__.py:174: UserWarning: Build succeeded, but resulted in non-empty logs: Build on &lt;pyopencl.Device 'Intel(R) Core(TM)2 CPU         T5600  @ 1.83GHz' at 0x1385a20&gt; succeeded, but said:
  
  Build started Kernel &lt;test&gt; was successfully vectorized Done.   warn("Build succeeded, but resulted in non-empty logs:\n"+message)
  [1]        # first value (ok)
  [67447488] # second value (wrong)
  
  Process finished with exit code 0
  </code></pre>