@@ -758,10 +758,19 @@ __select_backend(const execution::fpga_policy<_Factor, _KernelName>&, _Ranges&&.
758
758
}
759
759
#endif
760
760
761
- // Check the outer view type type to see if we can vectorize. Any non-contiguous inputs (e.g. reverse
762
- // views, permutation views, etc.) cannot be vectorized. If C++20 ranges are present, then we can
763
- // use the std::ranges::contiguous_range concept.
764
- template <typename _Rng, typename = void >
761
+ // TODO: At some point with C++20, we should implement this with concepts to more easily support the less common edge
762
+ // cases (e.g. a pipe over a counting iterator which is non-contiguous). For now, vectorization is primarily based on
763
+ // range contiguity with specialization for internal views and guard views over internal iterators.
764
+ // The following cases enable vectorization for a range:
765
+ // 1. With C++20 concepts, the range satisfies std::ranges::contiguous_range.
766
+ // 2. With C++17 nanoranges, the range satisfies __nanorange::nano::ranges::contiguous_range. Note that a view over
767
+ // a SYCL buffer satisfies this concept along with pipe views that maintain access contiguity.
768
+ // 3. The range is a guard view over an iterator with no global memory access: counting_iterator and discard_iterator
769
+ // 4. The range is one of our internal, vectorizable range types: drop_view_simple, take_view_simple, or
770
+ // transform_view_simple
771
+
772
+ // Base case: check contiguous range properties
773
+ template <typename _Rng>
765
774
struct __is_vectorizable_range
766
775
{
767
776
constexpr static bool value =
@@ -771,48 +780,18 @@ struct __is_vectorizable_range
771
780
__nanorange::nano::ranges::contiguous_range<_Rng>;
772
781
};
773
782
774
- // Guard view specializations
775
- // Counting iterator does not go through global memory but does not disable vectorization elsewhere.
783
+ // Basic guard view specializations - views which are not contiguous but do not interact with global memory
776
784
template <typename _Ip>
777
785
struct __is_vectorizable_range <oneapi::dpl::__ranges::guard_view<oneapi::dpl::counting_iterator<_Ip>>> : std::true_type
778
786
{
779
787
};
780
788
781
- // Discard iterator does not go through global memory but does not disable vectorization elsewhere.
782
789
template <>
783
790
struct __is_vectorizable_range <oneapi::dpl::__ranges::guard_view<oneapi::dpl::discard_iterator>> : std::true_type
784
791
{
785
792
};
786
793
787
- template <typename _Pointer>
788
- struct __is_vectorizable_range <oneapi::dpl::__ranges::guard_view<_Pointer>,
789
- std::enable_if_t <std::is_pointer_v<_Pointer>>> : std::true_type
790
- {
791
- };
792
-
793
- // For any non-pointer iterator over a guard_view, use contiguous iterator concepts
794
- template <typename _Iterator>
795
- struct __is_vectorizable_range <oneapi::dpl::__ranges::guard_view<_Iterator>,
796
- std::enable_if_t <!std::is_pointer_v<_Iterator>>>
797
- {
798
- constexpr static bool value =
799
- #if _ONEDPL_CPP20_RANGES_PRESENT && _ONEDPL_CPP20_CONCEPTS_PRESENT
800
- std::contiguous_iterator<_Iterator> ||
801
- #endif
802
- // std::vector iterators are not contiguous with nanorange so a separate check is necessary
803
- __nanorange::nano::contiguous_iterator<_Iterator> ||
804
- oneapi::dpl::__internal::__is_known_usm_vector_iter_v<_Iterator>;
805
- };
806
-
807
- // If all_view is passed, then we are processing a sycl::buffer directly which is contiguous and can
808
- // be used.
809
- template <typename _T, sycl::access::mode _AccMode, __dpl_sycl::__target _Target,
810
- sycl::access::placeholder _Placeholder>
811
- struct __is_vectorizable_range <oneapi::dpl::__ranges::all_view<_T, _AccMode, _Target, _Placeholder>> : std::true_type
812
- {
813
- };
814
-
815
- // Recursive view specializations - views which we need to search inwards to identify if it is vectorizable
794
+ // Recursive view specializations - internal views which we need to search inwards to identify if it is vectorizable
816
795
template <typename _Rng, typename _F>
817
796
struct __is_vectorizable_range <oneapi::dpl::__ranges::transform_view_simple<_Rng, _F>> : __is_vectorizable_range<_Rng>
818
797
{
0 commit comments