Do fp16 and fp64 aspects cover half * and double * data types?

[AlexeySachkov]

I've been looking at atomic_ref_add_sub_op_all_types_test_pointers.cpp CTS test failure on a device which does not support fp64 and realized that the spec is likely not clear enough about what exactly is guarded by fp16 and fp64 aspects.

atomic_ref tests for pointers unconditionally cover double * data type: atomic_ref_common.h:166. When implementation targets SPIR-V as intermediate representation for SYCL kernels, use of double * leads to emission of the following instructions sequence in SPIR-V module:

OpCapability Float64
%Double = OpTypeFloat 64
%DoblePtr = OpTypePointer %Double

Which makes it legal for a backend to discard such SPIR-V module if double is not supported by a target device. However, at SYCL level the app only uses double * and never dereferences that pointer, so there are no direct use of that data type.

The question is what was the intent of the spec with fp16 and fp64 aspects? Are they expected to also guard pointers to those optional data types?

I guess that from user experience point of view, pointers to optional data types should be allowed as long as they are not dereferenced on a device which doesn't support those types. Should this be clarified somehow in the SYCL 2020 spec?

Note: SPIR-V spec has special capability Float16Buffer capability, which allows pointers to half in a module, but doesn't allow their dereferencing. Unfortunately, there is no similar Float64Buffer capability.

[bader]

I guess that from user experience point of view, pointers to optional data types should be allowed as long as they are not dereferenced on a device which doesn't support those types.

I would say that would require typeless pointer representation in SPIR-V. I can imagine that double * might use different HW representation on FPGA device, even though CPU uses the same HW to represent pointers for any types.

Note: SPIR-V spec has special capability Float16Buffer capability, which allows pointers to half in a module, but doesn't allow their dereferencing. Unfortunately, there is no similar Float64Buffer capability.

Float16Buffer capability enables half dereferencing functionality via conversion to the wider type float. Float64 is the widest standard FP type, so it's unclear how someone use and implement Float64Buffer capability. How would you emulate double type operations using other (preferably non-optional) SPIR-V types?

[gmlueck]

My initial thought is similar to @bader, any use of the type double or sycl::half constitutes a use of the feature, even if it only declares a pointer type. We could perhaps add an exception for uses of double and sycl::half in unevaluated expressions, which means that something like sizeof(double) would not constitute a use.

[tomdeakin]

Please propose a PR for the CTS and the Spec.

[steffenlarsen]

CTS PR: KhronosGroup/SYCL-CTS#845

[tomdeakin]

CTS PR merged. Awaiting PR on spec.