vector_t WIP change

sycl/include/sycl/detail/builtins/helper_macros.hpp

@@ -192,12 +192,14 @@
  template <NUM_ARGS##_TYPENAME_TYPE> \
  detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>(NAME)( \
  NUM_ARGS##_TEMPLATE_TYPE_ARG) { \
+ /* vec(vector_t) is explicit, cannot rely on implicit conversion: */ \
+ using ret_type = detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>; \
  if constexpr (detail::is_marray_v<T0>) { \
- return detail::DELEGATOR(  \
+ return ret_type{detail::DELEGATOR( \
  [](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); }, \
- NUM_ARGS##_ARG);  \
+ NUM_ARGS##_ARG)}; \
  } else { \
- return __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG);  \
+ return ret_type{__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG)}; \
  } \
  }
 

sycl/include/sycl/detail/builtins/math_functions.inc

@@ -254,7 +254,10 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) {
  detail::NON_SCALAR_ENABLER<SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), \
  PtrTy> \
  NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) { \
- return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
+ using ret_ty = detail::NON_SCALAR_ENABLER< \
+ SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), PtrTy>; \
+ return ret_ty{ \
+ detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p)}; \
  }
 
 #if __SYCL_DEVICE_ONLY__

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/sub_group.hpp

@@ -81,7 +81,7 @@ vec<T, N> load(const multi_ptr<T, Space, DecorateAddress> src) {
  using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
  VecT Ret = __spirv_SubgroupBlockReadINTEL<VecT>(convertToBlockPtr(src));
 
- return sycl::bit_cast<typename vec<T, N>::vector_t>(Ret);
+ return vec<T, N>{sycl::bit_cast<typename vec<T, N>::vector_t>(Ret)};
 }
 
 template <typename T, access::address_space Space,

sycl/include/sycl/vector.hpp

@@ -1196,7 +1196,7 @@ template <typename Self,
 struct __SYCL_EBO VecConversionsMixin :
 #ifdef __SYCL_DEVICE_ONLY__
  public detail::ConversionOperatorMixin<
- Self, vector_t, ConversionOpType::conv_template,
+ Self, vector_t, ConversionOpType::conv_explicit,
  // if `vector_t` and `DataT` are the same, then the `operator DataT`
  // from the above is enough.
  !std::is_same_v<DataT, vector_t>>,
@@ -1581,13 +1581,12 @@ class __SYCL_EBO vec
 
 #ifdef __SYCL_DEVICE_ONLY__
  public:
- template <typename vector_t_,
- typename = std::enable_if_t<!std::is_same_v<vector_t_, DataT> &&
- std::is_same_v<vector_t_, vector_t>>>
+ template <typename vector_t_ = vector_t,
+ typename = std::enable_if_t<!std::is_same_v<vector_t_, DataT>>>
  // TODO: current draft would use non-template `vector_t` as an operand,
  // causing sycl::vec<sycl::half, N>{1} to go through different paths on
  // host/device, open question in the specification.
- constexpr vec(vector_t_ openclVector)
+ explicit vec(vector_t openclVector)
  // FIXME: Doesn't work when instantiated for 3-elements vectors,
  // indetermined padding can't be used to initialize constexpr std::array
  // storage.

sycl/test/basic_tests/types.cpp

@@ -202,8 +202,8 @@ void checkVariousVecUnaryOps() {
 
 int main() {
  // Test for creating constexpr expressions
- constexpr sycl::specialization_id<sycl::vec<sycl::half, 2>> id(1.0);
- constexpr sycl::marray<sycl::half, 2> MH(3);
+ constexpr sycl::specialization_id<sycl::vec<sycl::half, 2>> id(sycl::half{1.0});
+ constexpr sycl::marray<sycl::half, 2> MH(sycl::half{3});
  // Check the size and alignment of the SYCL vectors.
  checkVectors();
 

sycl/test/basic_tests/vectors/constexpr-constructor.cpp

@@ -6,10 +6,10 @@
 #include <cstdint>
 
 #define DEFINE_CONSTEXPR_VECTOR(name, type, size) \
- constexpr sycl::vec<type, size> name##_##size{0};
+ constexpr sycl::vec<type, size> name##_##size{(type){0}};
 
 #define DEFINE_CONSTEXPR_VECTOR_INIT_NON_ZERO(name, type, size, init) \
- constexpr sycl::vec<type, size> name##_##size{init};
+ constexpr sycl::vec<type, size> name##_##size{(type){init}};
 
 #define DEFINE_CONSTEXPR_VECTOR_FOR_TYPE(type) \
  DEFINE_CONSTEXPR_VECTOR(type, type, 1) \