sycl::vec's load method does not specify any aligment requirements.

[AD2605]

The description of sycl::vec's load method is as follows -

Loads the values at the address of ptr offset in elements of type DataT by NumElements * offset, into the components of this SYCL vec.

It does mention any alignment requirements imposed on the pointer being passed to the load method. Thus the following code snippet, would be in harmony with the sycl-spec, which would result in an error when run on certain backends -

void foo(float* ptr) {
    sycl::vec<float, 4> float_vector;
    char* a = reinterpret_cast<char*>(ptr);
   
   float_vector.load(0, a);
   ...
}

Similarly, the above example would hold true while using the store method as well
The specification should mention alignment impositions on the pointer being passed, to avoid allowing illegal

[gmlueck]

It's not clear to me what alignment restrictions there should be. I can see two possibilities:

1. Require ptr to be aligned according to section 4.14.2.6; or
2. Require ptr to be naturally aligned according to the DataT type.

I think your proposed spec change in #512 is attempting to specify option 2, correct?

If we want option 2, I'm not sure we need to change the spec at all. The parameter to load and store is a pointer to an object of type DataT, and the core C++ specification already requires that pointer to be naturally aligned.

BTW, I don't understand the code snippet in your comment above. The pointer a has the wrong type, but it is still correctly aligned for an object of type float.

[AD2605]

Hi Greg, Thanks for your reply.
I agree that the example is invalid. I missed the fact that the function signature accepts a multi_ptr of type DataT.

Nevertheless, I would still be of the opinion that the alignment of pointer that is being mentioned in the PR would be of value.

A bit of background, when I was using the load/store function, I was under the impression that using vec.load indeed would result in a vector load/store from the memory. However, when looking at the spec does not mention any such thing, and is thus implementation defined (That's also when I thought spec allowed for illegal behaviour, as in the example I provided, but as you mentioned, I stand corrected). If the expectation is indeed to do a vector load, then the pointer should be aligned to vector size.

Drawing a parallel to opencl's vloadn, it does mention the alignment requirement, and from that one can concur that it does not result in a vector load.

Adding that statement would set the expectation of the function to the user.

[AerialMantis]

SYCL WG call:

#512 has been merged now so we will close this, @AD2605 please re-open if you believe there is more to discuss here.