Restrict variadic vec ctor to N > 1

Author: aelovikov-intel

sycl/include/sycl/detail/spirv.hpp

@@ -871,11 +871,12 @@ EnableIfNativeShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
     return result;
   } else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
                            GroupT>) {
-    return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
-                                          convertToOpenCLType(x), LocalId);
+    return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
+        group_scope<GroupT>::value, convertToOpenCLType(x), LocalId));
   } else {
     // Subgroup.
-    return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId);
+    return convertFromOpenCLTypeFor<T>(
+        __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId));
   }
 #else
   if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -908,12 +909,12 @@ EnableIfNativeShuffle<T> ShuffleXor(GroupT g, T x, id<1> mask) {
     // general, and simple so we go with that.
     id<1> TargetLocalId = g.get_local_id() ^ mask;
     uint32_t TargetId = MapShuffleID(g, TargetLocalId);
-    return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
-                                          convertToOpenCLType(x), TargetId);
+    return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
+        group_scope<GroupT>::value, convertToOpenCLType(x), TargetId));
   } else {
     // Subgroup.
-    return __spirv_SubgroupShuffleXorINTEL(convertToOpenCLType(x),
-                                           static_cast<uint32_t>(mask.get(0)));
+    return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleXorINTEL(
+        convertToOpenCLType(x), static_cast<uint32_t>(mask.get(0))));
   }
 #else
   if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -956,12 +957,12 @@ EnableIfNativeShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
     if (TargetLocalId[0] + delta < g.get_local_linear_range())
       TargetLocalId[0] += delta;
     uint32_t TargetId = MapShuffleID(g, TargetLocalId);
-    return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
-                                          convertToOpenCLType(x), TargetId);
+    return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
+        group_scope<GroupT>::value, convertToOpenCLType(x), TargetId));
   } else {
     // Subgroup.
-    return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x),
-                                            convertToOpenCLType(x), delta);
+    return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleDownINTEL(
+        convertToOpenCLType(x), convertToOpenCLType(x), delta));
   }
 #else
   if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -1000,12 +1001,12 @@ EnableIfNativeShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
     if (TargetLocalId[0] >= delta)
       TargetLocalId[0] -= delta;
     uint32_t TargetId = MapShuffleID(g, TargetLocalId);
-    return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
-                                          convertToOpenCLType(x), TargetId);
+    return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
+        group_scope<GroupT>::value, convertToOpenCLType(x), TargetId));
   } else {
     // Subgroup.
-    return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x),
-                                          convertToOpenCLType(x), delta);
+    return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleUpINTEL(
+        convertToOpenCLType(x), convertToOpenCLType(x), delta));
   }
 #else
   if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<

sycl/include/sycl/stream.hpp

@@ -579,11 +579,11 @@ typename std::enable_if_t<(VecLength == 2 || VecLength == 4 || VecLength == 8 ||
                           unsigned>
 VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
          int Precision) {
-  unsigned Len =
-      VecToStr<T, VecLength / 2>(Vec.lo(), VecStr, Flags, Width, Precision);
+  unsigned Len = VecToStr<T, VecLength / 2>(vec<T, VecLength / 2>{Vec.lo()},
+                                            VecStr, Flags, Width, Precision);
   Len += append(VecStr + Len, VEC_ELEMENT_DELIMITER);
-  Len += VecToStr<T, VecLength / 2>(Vec.hi(), VecStr + Len, Flags, Width,
-                                    Precision);
+  Len += VecToStr<T, VecLength / 2>(vec<T, VecLength / 2>{Vec.hi()},
+                                    VecStr + Len, Flags, Width, Precision);
   return Len;
 }
 
@@ -593,7 +593,8 @@ VecToStr(const vec<T, VecLength> &Vec, char *VecStr, unsigned Flags, int Width,
          int Precision) {
   unsigned Len = VecToStr<T, 2>(Vec.lo(), VecStr, Flags, Width, Precision);
   Len += append(VecStr + Len, VEC_ELEMENT_DELIMITER);
-  Len += VecToStr<T, 1>(Vec.z(), VecStr + Len, Flags, Width, Precision);
+  Len +=
+      VecToStr<T, 1>(vec<T, 1>(Vec.z()), VecStr + Len, Flags, Width, Precision);
   return Len;
 }
 

sycl/include/sycl/vector.hpp

@@ -835,7 +835,9 @@ class __SYCL_EBO Swizzle
 
 #ifdef __SYCL_DEVICE_ONLY__
   operator vector_t() const {
-    return static_cast<vector_t>(static_cast<ResultVec>(*this));
+    // operator ResultVec() isn't available for single-element swizzle, create
+    // sycl::vec explicitly here.
+    return static_cast<vector_t>(ResultVec{this->Vec[Indexes]...});
   }
 #endif
 
@@ -1037,10 +1039,11 @@ class __SYCL_EBO vec :
 
   // Constructor from values of base type or vec of base type. Checks that
   // base types are match and that the NumElements == sum of lengths of args.
-  template <typename... argTN,
-            typename = std::enable_if_t<
-                ((AllowArgTypeInVariadicCtor<argTN> && ...)) &&
-                ((num_elements<argTN>() + ...)) == NumElements>>
+  template <
+      typename... argTN,
+      typename = std::enable_if_t<
+          (NumElements > 1 && ((AllowArgTypeInVariadicCtor<argTN> && ...)) &&
+           ((num_elements<argTN>() + ...)) == NumElements)>>
   constexpr vec(const argTN &...args)
       : vec{VecArgArrayCreator<DataT, argTN...>::Create(args...),
             std::make_index_sequence<NumElements>()} {}