atomic types (half2, bf162, float2, float4)

[jinz2014]

Is your feature request related to a problem? Please describe

error: static assertion failed due to requirement 'detail::IsValidAtomicRefType<sycl::vec<sycl::detail::half_impl::half, 2>>::value': Invalid atomic type. Valid types are int, unsigned int, long, unsigned long, long long, unsigned long long, sycl::half, float, double and pointer types

Describe the solution you would like

The float2 and float4 floating-point vector versions of atomicAdd() are only supported by devices of compute capability 9.x and higher. The atomicity of the float2 or float4 add operation is guaranteed separately for each of the two or four float elements; the entire float2 or float4 is not guaranteed to be atomic as a single 64-bit or 128-bit access.

The 16-bit __half floating-point version of atomicAdd() is only supported by devices of compute capability 7.x and higher.

The 16-bit __nv_bfloat16 floating-point version of atomicAdd() is only supported by devices of compute capability 8.x and higher.

The float2 and float4 floating-point vector versions of atomicAdd() are only supported by devices of compute capability 9.x and higher.

The float2 and float4 floating-point vector versions of atomicAdd() are only supported for global memory addresses.

Describe alternatives you have considered

No response

Additional context

No response

[Pennycook]

The atomicity of the float2 or float4 add operation is guaranteed separately for each of the two or four float elements; the entire float2 or float4 is not guaranteed to be atomic as a single 64-bit or 128-bit access.

This is not compatible with the semantics of std::atomic_ref, so it isn't compatible with the semantics of sycl::atomic_ref.

If you don't care whether the components of the vector are actually handled atomically, couldn't you create a separate sycl::atomic_ref for each component of the vector?

[AlexeySachkov]

Based on the comment from @Pennycook, I'm going to close the issue. atomic_ref by definition spans the whole data type passed to it. If you need a float2 version which does independent operations on individual components of that vector, then you can write your own helper function for that - it does not seem to be complicated enough to warrant an extension.