Skip to content

[ESIMD] Construct simd from simd_view using args deduction guide #4892

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 5 commits into from
Closed

[ESIMD] Construct simd from simd_view using args deduction guide #4892

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
c79f3f1
[ESIMD] Construct simd from simd_view using args deduction guide
v-klochkov Nov 3, 2021
f5ad03c
Merge remote-tracking branch 'intel_llvm/sycl' into esimd_simd_ctor_f…
v-klochkov Nov 4, 2021
71b6db0
Made the comment to the deduction guide a bit more descriptive
v-klochkov Nov 4, 2021
b7bc4dc
Add more comments to the new constructors
v-klochkov Nov 4, 2021
258a472
Merge remote-tracking branch 'intel_llvm/sycl' into esimd_simd_ctor_f…
v-klochkov Jan 21, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 13 additions & 0 deletions sycl/include/sycl/ext/intel/experimental/esimd/detail/simd_obj_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -213,6 +213,19 @@ class simd_obj_impl {
copy_from(acc, offset, Flags{});
}

/// Construct from \ref simd_view object \p View.
/// The constructed object has the same element type as the viewed simd.
/// and the size that is specified by the RegionT parameter of \p View.
///
/// \tparam ViewedSimdLength is the size in elements of the \ref simd
/// object being viewed by \p View.
/// \tparam RegionT is the region defining \p View.
/// \param View is the object from which the \ref simd is created.
template <int ViewedSimdLength, typename RegionT,
typename = std::enable_if_t<RegionT::length == N>>
simd_obj_impl(simd_view<simd<Ty, ViewedSimdLength>, RegionT> &View)
: M_data(View.read().M_data) {}

/// @}

// Load the object's value from array.
Expand Down
24 changes: 21 additions & 3 deletions sycl/include/sycl/ext/intel/experimental/esimd/simd.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,8 +51,8 @@ class simd : public detail::simd_obj_impl<
using raw_vector_type = typename base_type::raw_vector_type;
static constexpr int length = N;

// Implicit conversion constructor from another simd object of the same
// length.
/// Implicit conversion constructor from another simd object of the same
/// length.
template <typename SimdT,
class = std::enable_if_t<__SEIEED::is_simd_type_v<SimdT> &&
(length == SimdT::length)>>
Expand All @@ -62,13 +62,26 @@ class simd : public detail::simd_obj_impl<
__esimd_dbg_print(simd(const SimdT &RHS));
}

// Broadcast constructor with conversion.
/// Broadcast constructor with conversion.
template <typename T1,
class = std::enable_if_t<detail::is_valid_simd_elem_type_v<T1>>>
simd(T1 Val) : base_type(Val) {
__esimd_dbg_print(simd(T1 Val));
}

/// Construct \ref simd from \ref simd_view object \p View.
/// The constructed \ref simd object has the same element type as the viewed
/// simd and the size that is specified by the RegionT parameter of \p View.
///
/// \tparam ViewedSimdLength is the size in elements of the \ref simd
/// object being viewed by \p View.
/// \tparam RegionT is the region defining the \p View.
/// \param View is the object from which the \ref simd is created.
template <int ViewedSimdLength, typename RegionT,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

ditto

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed in b7bc4dc

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks!
Seems like this is almost complete copy-paste of the simd_obj_impl variant. I missed that in the first review, unfortunately. Can duplication be somehow avoided? E.g. this simd variant just refers to the simd_obj_impl variant?

One more thing that I missed is simd_mask - it should also have this kind of constructor. I wonder if it is enough just to define that constructor in a generic way so that it is automatically inherited from simd_obj_impl by both simd and simd_mask w/o the need to define this constructor explicitly in them?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Technically it can be deleted, I'll do that.

Need to add though that the deduction guide needs to be defined for 'simd' class, not for 'simd_obj_impl'.
Having something like this will now work for the assignment: "simd_view SV = V; simd V1 = SV":

// Deduction guide for 'simd_obj_impl' constructor from a 'simd_view' object.
// It helps to infer the template parameter 'N' of the class 'simd_obj_impl'.
template <typename T, int ViewedSimdLength, typename RegionT>
simd_obj_impl(simd_view<simd<T, ViewedSimdLength>, RegionT>) -> simd_obj_impl<T, RegionT::length, simd<T, RegionT::length>>;

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

One more note: the mentioned 'simd' constructor can be deleted only because of this statement in simd class:

using base_type::base_type;

The same for 'simd -> simd_view' is not doable because 'simd_view' does not have such 'using' statement and the following ctor cannot be removed:
https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/intel/experimental/esimd/simd_view.hpp#L76

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Regarding the 'simd_mask' and doing the arguments deduction for it...
Currently, 'simd_mask' is declared as an alias template, not as a template class and thus the template args deduction guides do not work for it.

simd_mask<16> m16 = 0;
auto mask_view8 = m16.select<8, 1>(0);

detail::simd_mask_impl m8 = mask_view8; // Works well with some changes in headers
simd_mask m8 = mask_view8; // Does not work. The error is: "error: alias template 'simd_mask' requires template arguments; argument deduction only allowed for class templates"

I don't know why simd_mask is not defined as a class instead of an alias to a class in detail namespace.
Anyway, tackling this problem for simd_mask deserves a separate PR, IMO.

typename = std::enable_if_t<RegionT::length == N>>
simd(simd_view<simd<element_type, ViewedSimdLength>, RegionT> &View)
: base_type(View) {}

/// Type conversion for simd<T, 1> into T.
template <class To, class T = simd,
class = sycl::detail::enable_if_t<
Expand Down Expand Up @@ -116,6 +129,11 @@ class simd : public detail::simd_obj_impl<
#undef __ESIMD_DEF_SIMD_ARITH_UNARY_OP
};

// Deduction guide for 'simd' constructor from a 'simd_view' object. It helps
// to infer the template parameter 'N' of the class 'simd'.
template <typename T, int ViewedSimdLength, typename RegionT>
simd(simd_view<simd<T, ViewedSimdLength>, RegionT>) -> simd<T, RegionT::length>;

/// Covert from a simd object with element type \c From to a simd object with
/// element type \c To.
template <typename To, typename From, int N>
Expand Down
10 changes: 10 additions & 0 deletions sycl/test/esimd/simd_view.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,6 +124,16 @@ SYCL_ESIMD_FUNCTION void test_simd_view_from_vector() {
simd_view sv1b(v1);
}

// test deducing arguments during simd construction from simd_view.
SYCL_ESIMD_FUNCTION void test_simd_from_simd_view() {
simd<int, 16> v16 = 0;
auto view8 = v16.select<8,1>(0);
simd v8 = view8;

auto view1 = v16.select<1,1>(0);
simd v1 = view1;
}

// move constructor transfers the same view of the underlying data.
SYCL_ESIMD_FUNCTION void test_simd_view_move_ctor() {
simd<int, 16> v0 = 1;
Expand Down