Archives Discussions

[firespot]

Hi,

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?

thanks!

[dipak · Accepted solution]

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 vloadto read vectors from scalar pointer.

Regards

[jammyamerica]

Helpfull:)