argument has type 'volatile sycl::float4' (aka 'volatile vec<float, 4>'), but method is not marked volatile

[zjin-lcf]

I converted the cuda code (https://github.com/zjin-lcf/HeCBench/tree/master/bh-cuda) to sycl with Intel DPCT, and then compiled the sycl codes.

The error message is shown below. Could the method be marked volatile ? Thanks.

/path/to/build/bin/../include/sycl/CL/sycl/swizzles.def:79:50: note: candidate function template not viable: 'this' argument has type 'volatile sycl::float4' (aka 'volatile vec<float, 4>'), but method is not marked volatile
__SYCL_SCALAR_ACCESS(N == 2 || N == 3 || N == 4, y, 1)
^
/path/to/build/bin/../include/sycl/CL/sycl/swizzles.def:79:50: note: candidate function template not viable: 'this' argument has type 'volatile sycl::float4' (aka 'volatile vec<float, 4>'), but method is not marked volatile
main.dp.cpp:416:46: error: no matching member function for call to 'z'
const float chz = posMassd[ch].z();

[bader]

@intel/dpcpp-specification-reviewers, should __swizzled_vec__ XYZW_ACCESS() const; vec member has volatile qualifier?
https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_vec_interface

[bader]

volatile works with std::experimetal::simd: https://godbolt.org/z/55PqYvGvj

[zjin-lcf]

Will experimental 'simd' be merged into sycl::vec ? Thanks.

[bader]

    volatile native_simd<float> c{};
    float f = c[0];

Actually I missed the critical part - calling a member function of the class. simd also gives an error.

Will experimental 'simd' be merged into sycl::vec ? Thanks.

I provided an example with simd as I think it represents the concept sycl::vec is supposed to represent more closely than std::array.

[gmlueck]

Getting back to the original question:

should swizzled_vec XYZW_ACCESS() const; vec member has volatile qualifier?

It does not seem like this makes sense. There is no precedent in C++ standard containers to add volatile to their member functions. In addition, it seems unlikely that it's really correct to use volatile in SYCL device code. This is the meaning of volatile from cppreference:

Every access (read or write operation, member function call, etc.) made through a glvalue expression of volatile-qualified type is treated as a visible side-effect for the purposes of optimization (that is, within a single thread of execution, volatile accesses cannot be optimized out or reordered with another visible side effect that is sequenced-before or sequenced-after the volatile access. This makes volatile objects suitable for communication with a signal handler, but not with another thread of execution, see std::memory_order).

Note my emphasized text. Since volatile is not a suitable way to communicate with another thread, my guess is that the original CUDA program is not technically correct. Perhaps using volatile works in some cases, but it is not guaranteed to work by the C++ specification.

If you want to take your chances, you could use const_cast to cast away the volatile qualifier before calling the member function. However, there is no guarantee that this will synchronize correctly with another thread, and this might be the reason the converted version of the application hangs. A better solution would be to understand the inter-thread synchronization in your application and use the SYCL "atomic" functions instead.

[zjin-lcf]

I.4.3.3. Volatile Qualifier
The compiler is free to optimize reads and writes to global or shared memory (for example, by caching global reads into registers or L1 cache) as long as it respects the memory ordering semantics of memory fence functions (Memory Fence Functions) and memory visibility semantics of synchronization functions (Synchronization Functions).

These optimizations can be disabled using the volatile keyword: If a variable located in global or shared memory is declared as volatile, the compiler assumes that its value can be changed or used at any time by another thread and therefore any reference to this variable compiles to an actual memory read or write instruction.

[zjin-lcf]

Thank you for your example and description. My understanding is that the cuda example requires "volatile" for "float4".
Running the DPCT version hangs.