[SYCL] Make vec conversion operator to scalar non-template (#14668)

aelovikov-intel · web-flow · commit c2ebf84fd7ff · 2024-07-22T13:57:54.000-07:00
Template vs. non-template conversion operator behave differently for
```
  sycl::vec&lt;int, 1&gt; v;
  std::ignore = static_cast&lt;bool&gt;(v);
```

where non-template path selects the conversion operator and then uses
implicit `int-&gt;bool` conversion while the example fails to compile when
operator is a template.

This is important because the spec requires `vec::x()` and other
swizzles to return a swizzle class and not a `const DataT &amp;` that it
currently does. However, if we'd fix the swizzle without addressing this
issues first the code like
```
  sycl::vec&lt;int, 8&gt; v = {...};
  if (v.x() == 42)
    ...
```
would fail to compile. Also, the spec is explicit that this conversion operator
must not be `template`:
 
&gt; The SYCL vec instance shall be implicitly convertible to the same data types,
&gt; to which DataT is implicitly convertible. Note that conversion operator shall 
&gt; not be templated to allow standard conversion sequence for implicit conversion.
diff --git a/sycl/include/sycl/aliases.hpp b/sycl/include/sycl/aliases.hpp
@@ -15,7 +15,7 @@
 
 namespace sycl {
 inline namespace _V1 {
-template <typename T, int N> class vec;
+template <typename T, int N> class __SYCL_EBO vec;
 namespace detail::half_impl {
 class half;
 } // namespace detail::half_impl
diff --git a/sycl/include/sycl/detail/common.hpp b/sycl/include/sycl/detail/common.hpp
@@ -167,14 +167,6 @@ class __SYCL_EXPORT tls_code_loc_t {
 
 #include <sycl/exception.hpp>
 
-// Helper for enabling empty-base optimizations on MSVC.
-// TODO: Remove this when MSVC has this optimization enabled by default.
-#ifdef _MSC_VER
-#define __SYCL_EBO __declspec(empty_bases)
-#else
-#define __SYCL_EBO
-#endif
-
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
diff --git a/sycl/include/sycl/detail/defines_elementary.hpp b/sycl/include/sycl/detail/defines_elementary.hpp
@@ -35,6 +35,15 @@
 #endif
 #endif
 
+// Helper for enabling empty-base optimizations on MSVC.
+// TODO: Remove this when MSVC has this optimization enabled by default.
+#ifdef _MSC_VER
+#define __SYCL_EBO __declspec(empty_bases)
+#else
+#define __SYCL_EBO
+#endif
+
+
 #ifndef __SYCL_ID_QUERIES_FIT_IN_INT__
 #define __SYCL_ID_QUERIES_FIT_IN_INT__ 0
 #endif
diff --git a/sycl/include/sycl/detail/generic_type_lists.hpp b/sycl/include/sycl/detail/generic_type_lists.hpp
@@ -20,7 +20,7 @@
 // Forward declarations
 namespace sycl {
 inline namespace _V1 {
-template <typename T, int N> class vec;
+template <typename T, int N> class __SYCL_EBO vec;
 template <typename Type, std::size_t NumElements> class marray;
 
 namespace detail {
diff --git a/sycl/include/sycl/detail/vector_arith.hpp b/sycl/include/sycl/detail/vector_arith.hpp
@@ -21,7 +21,7 @@
 namespace sycl {
 inline namespace _V1 {
 
-template <typename DataT, int NumElem> class vec;
+template <typename DataT, int NumElem> class __SYCL_EBO vec;
 
 namespace detail {
 
@@ -287,7 +287,7 @@ class vec_arith : public vec_arith_common<DataT, NumElements> {
   __SYCL_BINOP(<<, <<=, true, (!detail::is_vgenfloat_v<DataT>))
 
   // friends
-  template <typename T1, int T2> friend class vec;
+  template <typename T1, int T2> friend class __SYCL_EBO vec;
 }; // class vec_arith<>
 
 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
@@ -340,7 +340,7 @@ class vec_arith<std::byte, NumElements>
   __SYCL_BINOP(^, ^=, false, true)
 
   // friends
-  template <typename T1, int T2> friend class vec;
+  template <typename T1, int T2> friend class __SYCL_EBO vec;
 };
 #endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 
@@ -383,7 +383,7 @@ template <typename DataT, int NumElements> class vec_arith_common {
 #endif
 
   // friends
-  template <typename T1, int T2> friend class vec;
+  template <typename T1, int T2> friend class __SYCL_EBO vec;
 };
 
 #undef __SYCL_BINOP
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
@@ -27,7 +27,7 @@ inline namespace _V1 {
 template <typename T, access::address_space Space,
           access::decorated DecorateAddress>
 class multi_ptr;
-template <typename Type, int NumElements> class vec;
+template <typename Type, int NumElements> class __SYCL_EBO vec;
 namespace detail {
 
 namespace sub_group {
diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp
@@ -97,13 +97,38 @@ template <typename T> class GetOp {
   DataT operator()(DataT, DataT) { return (DataT)0; }
 };
 
+// Templated vs. non-templated conversion operator behaves differently when two
+// conversions are needed as in the case below:
+//
+//   sycl::vec<int, 1> v;
+//   std::ignore = static_cast<bool>(v);
+//
+// Make sure the snippet above compiles. That is important because
+//
+//   sycl::vec<int, 2> v;
+//   if (v.x() == 42)
+//     ...
+//
+// must go throw `v.x()` returning a swizzle, then its `operator==` returning
+// vec<int, 1> and we want that code to compile.
+template <typename Vec, typename T, int N, typename = void>
+struct ScalarConversionOperatorMixIn {};
+
+template <typename Vec, typename T, int N>
+struct ScalarConversionOperatorMixIn<Vec, T, N, std::enable_if_t<N == 1>> {
+  operator T() const { return (*static_cast<const Vec *>(this))[0]; }
+};
+
 } // namespace detail
 
 ///////////////////////// class sycl::vec /////////////////////////
 // Provides a cross-platform vector class template that works efficiently on
 // SYCL devices as well as in host C++ code.
 template <typename DataT, int NumElements>
-class vec : public detail::vec_arith<DataT, NumElements> {
+class __SYCL_EBO vec
+    : public detail::vec_arith<DataT, NumElements>,
+      public detail::ScalarConversionOperatorMixIn<vec<DataT, NumElements>,
+                                                   DataT, NumElements> {
 
   static_assert(NumElements == 1 || NumElements == 2 || NumElements == 3 ||
                     NumElements == 4 || NumElements == 8 || NumElements == 16,
@@ -373,12 +398,6 @@ class vec : public detail::vec_arith<DataT, NumElements> {
   operator vector_t() const { return sycl::bit_cast<vector_t>(m_Data); }
 #endif // __SYCL_DEVICE_ONLY__
 
-  // Available only when: NumElements == 1
-  template <int N = NumElements>
-  operator typename std::enable_if_t<N == 1, DataT>() const {
-    return m_Data[0];
-  }
-
   __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead")
   static constexpr size_t get_count() { return size(); }
   static constexpr size_t size() noexcept { return NumElements; }
@@ -605,7 +624,7 @@ class vec : public detail::vec_arith<DataT, NumElements> {
   template <typename T1, typename T2, typename T3, template <typename> class T4,
             int... T5>
   friend class detail::SwizzleOp;
-  template <typename T1, int T2> friend class vec;
+  template <typename T1, int T2> friend class __SYCL_EBO vec;
   // To allow arithmetic operators access private members of vec.
   template <typename T1, int T2> friend class detail::vec_arith;
   template <typename T1, int T2> friend class detail::vec_arith_common;
diff --git a/sycl/test/abi/layout_vec.cpp b/sycl/test/abi/layout_vec.cpp
@@ -10,22 +10,20 @@
 
 SYCL_EXTERNAL void foo(sycl::vec<int, 4>) {}
 
-// CHECK: 0 | class sycl::vec<int, 4>
-// CHECK: 0 |   class sycl::detail::vec_arith<int, 4> (base) (empty)
-// CHECK-NEXT: 0 |     class sycl::detail::vec_arith_common<int, 4> (base) (empty)
-// CHECK-NEXT: 0 |   struct std::array<int, 4> m_Data
+// CHECK:      0 | class sycl::vec<int, 4>
+// ignore empty base classes
+// CHECK:      0 |   struct std::array<int, 4> m_Data
 // CHECK-NEXT: 0 |     typename _AT_Type::_Type _M_elems
-// CHECK-NEXT: | [sizeof=16, dsize=16, align=16,
-// CHECK-NEXT: |  nvsize=16, nvalign=16]
+// CHECK-NEXT:   | [sizeof=16, dsize=16, align=16,
+// CHECK-NEXT:   |  nvsize=16, nvalign=16]
 
 //--------------------------------------
 
 SYCL_EXTERNAL void foo(sycl::vec<bool, 16>) {}
 
-// CHECK: 0 | class sycl::vec<_Bool, 16>
-// CHECK: 0 |   class sycl::detail::vec_arith<_Bool, 16> (base) (empty)
-// CHECK-NEXT: 0 |     class sycl::detail::vec_arith_common<_Bool, 16> (base) (empty)
-// CHECK-NEXT: 0 |   struct std::array<_Bool, 16> m_Data
+// CHECK:      0 | class sycl::vec<_Bool, 16>
+// ignore empty base classes
+// CHECK:      0 |   struct std::array<_Bool, 16> m_Data
 // CHECK-NEXT: 0 |     typename _AT_Type::_Type _M_elems
-// CHECK-NEXT: | [sizeof=16, dsize=16, align=16,
-// CHECK-NEXT: |  nvsize=16, nvalign=16]
+// CHECK-NEXT:   | [sizeof=16, dsize=16, align=16,
+// CHECK-NEXT:   |  nvsize=16, nvalign=16]
diff --git a/sycl/test/basic_tests/vectors/size_one_checks.cpp b/sycl/test/basic_tests/vectors/size_one_checks.cpp
@@ -0,0 +1,14 @@
+// RUN: %clangxx -fsycl %s -o %t_default.out
+// RUN: %t_default.out
+// RUN: %if preview-breaking-changes-supported %{ %clangxx -fsycl -fpreview-breaking-changes %s -o %t_vec.out %}
+// RUN: %if preview-breaking-changes-supported %{ %t_vec.out %}
+
+#include <sycl/types.hpp>
+
+int main() {
+  sycl::vec<int, 1> v1{42};
+  sycl::vec<int, 1> v2{0};
+  assert(static_cast<bool>(v1) == true);
+  assert(static_cast<bool>(v2) == false);
+  return 0;
+}
diff --git a/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp b/sycl/test/check_device_code/vector/vector_bf16_builtins.cpp
diff --git a/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp b/sycl/test/check_device_code/vector/vector_convert_bfloat.cpp
diff --git a/sycl/test/check_device_code/vector/vector_math_ops.cpp b/sycl/test/check_device_code/vector/vector_math_ops.cpp
