2 Replies Latest reply on Aug 13, 2014 4:12 AM by jammyamerica

    alignment of vector types




      Consider an array of ints to be read using vector types:


      __kernel void someKernel(__global int const * values)


      int8 A = *((__global int8 const *) (values)); // no issues I suppose

      int8 B = *((__global int8 const *) (values+ 1)); // alignment catch?



      I guess the first access is always ok. Is the second also always ok, or can one be hit by alignment issues, e.g. if an int8 vector type is always assumed to be properly aligned on a 4*8 = 32 bytes boundary at access?



        • Re: alignment of vector types

          As per the OpenCL spec 1.2:

          6.2.5 Pointer Casting

          Pointers to old and new types may be cast back and forth to each other. Casting a pointer to a

          new type represents an unchecked assertion that the address is correctly aligned. The developer

          will also need to know the endianness of the OpenCL device and the endianness of the data to

          determine how the scalar and vector data elements are stored in memory.

          As int8* needs to be 8*4 = 32 bytes aligned, 2nd statement will create an alignment issue.


          Instead of type casting, its safe to use vload[n]to read vectors from scalar pointer.