From a57f1a78be334ef3e53e382ff170ab170e2e54e1 Mon Sep 17 00:00:00 2001 From: Greg Lueck Date: Mon, 2 Dec 2024 16:37:09 -0500 Subject: [PATCH] Remove underspecified vector_t This is change 9 of 9 that fix problems with the specification of the `vec` class. An implementation that follows the existing specification would not accept common code patterns and would not pass the CTS. None of the existing implementations actually follow the existing specification. This change removes the `vector_t` type alias from the `vec` class and also removes the conversions to and from this type. The intention of this type was to provide a conversion between `vec` and a compiler-specific builtin vector type. However, these conversion also lead to ambiguity errors in common code patterns such as: ``` sycl::vec h1; sycl::half h; h = h1; ``` Prior to this change, the last line could result in an ambiguity error from the compiler between: * Convert `vec` -> `half`, call `half::operator=(const half&)` * Convert `vec` -> `_Float16`, convert `_Float16` -> `float`, call `half::operator=(float)` What's worse is that the exact type alias of `vector_t` is implementation-defined, so these ambiguity errors could be different for each SYCL implementation. During the WG meeting, we learned that the two major SYCL implementations (DPC++ and AdaptiveCpp) define `vector_t` differently even though they are both based on clang. Given this lack of commonality, the WG decided to remove `vector_t` entirely from the spec. Implementations can decide for themselves whether to provide a compiler-specific conversion like this. These changes correspond to slides 29 - 31 of the presentation that was discussed in the WG meetings. --- adoc/chapters/opencl_backend.adoc | 4 +- adoc/chapters/programming_interface.adoc | 55 ++---------------------- adoc/headers/vec.h | 10 ----- 3 files changed, 6 insertions(+), 63 deletions(-) diff --git a/adoc/chapters/opencl_backend.adoc b/adoc/chapters/opencl_backend.adoc index 37bc993c..c7efb008 100644 --- a/adoc/chapters/opencl_backend.adoc +++ b/adoc/chapters/opencl_backend.adoc @@ -1063,8 +1063,8 @@ OpenCL C functions that are linked with, using either offline or online compilation, must be declared as [code]#extern "C"# function declarations. The function parameters of these function declarations must be defined as the OpenCL C interoperability aliases; [code]#pointer# of the [code]#multi_ptr# -class template, [code]#vector_t# of the [code]#vec# class template and scalar -data type aliases described in <>. +class template, and scalar data type aliases described in +<>. // \include{opencl_extensions} // %%%%%%%%%%%%%%%%%%%%%%%%%%%% begin opencl_extensions %%%%%%%%%%%%%%%%%%%%%%%%%%%% diff --git a/adoc/chapters/programming_interface.adoc b/adoc/chapters/programming_interface.adoc index 5f2b2f9c..e268e3cf 100644 --- a/adoc/chapters/programming_interface.adoc +++ b/adoc/chapters/programming_interface.adoc @@ -17763,18 +17763,10 @@ Any other value shall produce a compilation failure. The element type parameter, _DataT_, must be one of the basic scalar types supported in device code. -The SYCL [code]#vec# class template provides interoperability with the -underlying vector type defined by [code]#vector_t# which is available only when -compiled for the device. -The SYCL [code]#vec# class can be constructed from an instance of -[code]#vector_t# and can implicitly convert to an instance of [code]#vector_t# -in order to support interoperability with native <> functions from a -SYCL kernel function. - -An instance of the SYCL [code]#vec# class template can also be implicitly -converted to an instance of the data type when the number of elements is -[code]#1# in order to allow single element vectors and scalars to be convertible -with each other. +An instance of the SYCL [code]#vec# class template can be implicitly converted +to an instance of the data type when the number of elements is [code]#1# in +order to allow single element vectors and scalars to be convertible with each +other. ==== Vec interface @@ -17826,15 +17818,6 @@ constexpr vec(const vec& rhs) ---- a@ Construct a vector of element type [code]#DataT# and number of elements [code]#NumElements# by copy from another similar vector. -a@ -[source] ----- -vec(vector_t nativeVector) ----- - a@ Available only when: compiled for the device. - -Constructs a SYCL [code]#vec# instance from an instance of the underlying backend-native vector type defined by [code]#vector_t#. - |==== @@ -17844,16 +17827,6 @@ Constructs a SYCL [code]#vec# instance from an instance of the underlying backen [width="100%",options="header",separator="@",cols="65%,35%"] |==== @ Member function @ Description -a@ -[source] ----- -operator vector_t() const ----- - a@ Available only when: compiled for the device. - -Converts this SYCL [code]#vec# instance to the underlying backend-native vector type -defined by [code]#vector_t#. - a@ [source] ---- @@ -18598,22 +18571,6 @@ using value_type = DataT Each of these type aliases tells the type of an element in the underlying [code]#vec#. -''' - -[frame=all,grid=none,separator="@"] -|==== -a@ -[source] ----- -#ifdef __SYCL_DEVICE_ONLY__ -using vector_t = /*unspecified*/ -#endif ----- -|==== -This type alias is available only in device code. -The [code]#vector_t# type represents a native vector type of [code]#NumElements# -elements where each element's type is [code]#DataT#. - [[sec:constructors.swizzled-vec]] ===== Constructors for the swizzled vector class templates @@ -18656,10 +18613,6 @@ The destructors have no visible effect. a@ [source] ---- -#ifdef __SYCL_DEVICE_ONLY__ -operator vector_t() const -#endif - operator DataT() const static constexpr size_t byte_size() noexcept static constexpr size_t size() noexcept diff --git a/adoc/headers/vec.h b/adoc/headers/vec.h index 854294e2..8961289d 100644 --- a/adoc/headers/vec.h +++ b/adoc/headers/vec.h @@ -37,10 +37,6 @@ template class vec { using element_type = DataT; using value_type = DataT; -#ifdef __SYCL_DEVICE_ONLY__ - using vector_t = __unspecified__; -#endif - vec(); explicit constexpr vec(const DataT& arg); @@ -49,12 +45,6 @@ template class vec { constexpr vec(const vec& rhs); -#ifdef __SYCL_DEVICE_ONLY__ - vec(vector_t nativeVector); - - operator vector_t() const; -#endif - // Available only when: NumElements == 1 operator DataT() const;