[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513

MikeDvorskiy · 2024-04-18T09:05:27Z

[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types.

After offline discussed internally in oneDPL team, we decide to keep sycl::is_device_copyable trait specialization next to each oneDPL type definition.
Also we have to support the previous oneAPI compiler version via special macro
Also we have to support a case when SYCL backend is not active.

This is a prototype for 3 different oneDPL types: transform_reduce, __brick_copy_n and zip_iterator
(The total number of such oneDPL types is about 50.)
Also the prototype introduces a common code stuff to define the is_device_copyable trait specializations.

Based on GoDBolt prototype: https://godbolt.org/z/qx8nTaKsK

UPDATE!!!
The prototype was modified: via a special trait in each device copyable oneDPL type

…or transform_reduce

…or brick_copy_n and zip_iterator

…EVICE_COPYABLE definition into the global namespace

danhoeflinger

After our investigation and offline discussions, I agree that this is the best we can do given the context:

non-type params
variable # of types
some arbitrary subset of template params
supporting different version of llvm compiler (breaking change)

This fixes a few technical issues with the existing system, and also mitigates the maintenance burden which was caused by separating the device copyable specializations from the types themselves.

We will of course need to remove the old device copyable system / header in the eventual PR.

danhoeflinger · 2024-04-18T11:43:05Z

include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h

+
+template <typename _ExecutionPolicy, ::std::uint8_t __iters_per_work_item, typename _Operation1, typename _Operation2,
+          typename _Tp, typename _Commutative>
+struct _ONEDPL_IS_DEVICE_COPYABLE(oneapi::dpl::unseq_backend::transform_reduce, _ExecutionPolicy, __iters_per_work_item,


For other reviewers:
This case is interesting because of the non-type template parameter __iters_per_work_item. This causes problems for matching against the previous version of the oneDPL device copyable macro with a variadic template typename pack.

danhoeflinger · 2024-04-18T11:57:00Z

include/oneapi/dpl/pstl/algorithm_fwd.h

+} //__internal
+} //dpl
+} //oneapi


I wish there was a way to cleanly "exit" scope temporarily, to define within ::sycl, but I don't know of a better way than this.

danhoeflinger · 2024-04-18T12:05:43Z

include/oneapi/dpl/pstl/utils.h

@@ -48,6 +48,30 @@ namespace dpl
 namespace __internal
 {

+#if _ONEDPL_BACKEND_SYCL
+
+template <typename... _Ts>


In the full PR, we need comments for how this is to be used:
"__is_types_device_copyable is used to define which template parameters are required to be sycl device copyable to make the type being specialized device copyable. Generally, this means the template arguments which are used as member variables of the type being specialized."

danhoeflinger · 2024-04-18T12:15:17Z

include/oneapi/dpl/pstl/utils.h

+template <typename... _Ts>
+struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};
+
+#if __INTEL_LLVM_COMPILER && (__INTEL_LLVM_COMPILER < 20240100)


Proposed comment:

// _ONEDPL_IS_DEVICE_COPYABLE is used to declare the type we are specializing to be device copyable, with the // template type name listed first, and then each template argument listed next. The following is an example for // a structure A with two template parameters, where only the second template parameter type is used as a // member variable: // // template <typename T1, typename T2> struct A{ T2 b; }; // // template<typename T1, typename T2> // struct _ONEDPL_IS_DEVICE_COPYABLE(A, T1, T2): // oneapi::dpl::__internal::__is_types_device_copyable<T2> {}; //

danhoeflinger · 2024-04-18T12:16:06Z

include/oneapi/dpl/pstl/utils.h

+#   define _ONEDPL_IS_DEVICE_COPYABLE(TYPE, ...) sycl::is_device_copyable<TYPE<__VA_ARGS__>>
+#endif //__INTEL_LLVM_COMPILER && (__INTEL_LLVM_COMPILER < 20240100)
+
+#else


Suggested change

#else

#else // _ONEDPL_BACKEND_SYCL

danhoeflinger · 2024-04-18T12:17:07Z

include/oneapi/dpl/pstl/utils.h

+#if _ONEDPL_BACKEND_SYCL
+
+template <typename... _Ts>
+struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};


Suggested change

struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};

struct __are_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};

danhoeflinger · 2024-04-18T15:58:10Z

include/oneapi/dpl/pstl/algorithm_impl.h

+
+    _ONEDPL_IS_DEVICE_COPYABLE()


I think this is unnecessary, and will be so whenever there are no types to provide to this macro.
The type is already "trivially copyable", so it is inherently "device copyable".

Additionally, this is in the non-hetero impl so this will always result in a noop.

I want to walk that back slightly. It could be conceivable that there could exist a type which we just want to always claim is device copyable, with no arguments here.

However, most (if not all) cases will be either "trivially copyable", or have some fixed type we would want to send in to this macro. For example, if a type always had a std::tuple<int,int> member field unrelated to template arguments, we would need to "mark" it as "device copyable". We could use _ONEDPL_IS_DEVICE_COPYABLE(), but it would perhaps be more clear to instead use _ONEDPL_IS_DEVICE_COPYABLE(std::tuple<int, int>) to better document the reason it is required in the first place.

danhoeflinger · 2024-04-18T15:59:59Z

include/oneapi/dpl/pstl/hetero/algorithm_impl_hetero.h

@@ -375,6 +375,8 @@ struct __brick_copy_n<__hetero_tag<_BackendTag>, _ExecutionPolicy>
    {
        __target = ::std::forward<_SourceT>(__source);
    }
+
+    _ONEDPL_IS_DEVICE_COPYABLE()


Unnecessary, as mentioned above. I think __brick_fill_n may be the one you intended to add to.

…ARIANT 2: via a special trait in each device copyable oneDPL type

MikeDvorskiy · 2024-05-07T15:13:24Z

Close the draft because the other solution (right solution) was merged. See #1521 for details.

[oneDPL][hetero] + __is_types_device_copyable trait, specialization f…

1597a35

…or transform_reduce

MikeDvorskiy requested review from akukanov, rarutyun and danhoeflinger April 18, 2024 09:05

MikeDvorskiy marked this pull request as draft April 18, 2024 09:05

[oneDPL][hetero] + __is_types_device_copyable trait, specialization f…

3be871d

…or brick_copy_n and zip_iterator

MikeDvorskiy force-pushed the dev/mdvorski/dev_copyable_traits branch from 0092c9b to 3be871d Compare April 18, 2024 09:07

MikeDvorskiy mentioned this pull request Apr 18, 2024

Improvements / fixes to device copyable specializations #1505

Closed

[oneDPL][hetero] device copyable trait definition: bring _ONEDPL_IS_D…

0ed2d1b

…EVICE_COPYABLE definition into the global namespace

MikeDvorskiy force-pushed the dev/mdvorski/dev_copyable_traits branch from fa945ba to 0ed2d1b Compare April 18, 2024 10:56

danhoeflinger reviewed Apr 18, 2024

View reviewed changes

danhoeflinger linked an issue Apr 18, 2024 that may be closed by this pull request

Improvements / fixes to device copyable specializations #1505

Closed

MikeDvorskiy force-pushed the dev/mdvorski/dev_copyable_traits branch from d5b1d78 to 0539960 Compare April 18, 2024 15:52

danhoeflinger reviewed Apr 18, 2024

View reviewed changes

[oneDPL][hetero] device copyable trait definitions: implementation, V…

f4c12b1

…ARIANT 2: via a special trait in each device copyable oneDPL type

MikeDvorskiy force-pushed the dev/mdvorski/dev_copyable_traits branch from 0539960 to f4c12b1 Compare April 19, 2024 10:13

This was referenced Apr 19, 2024

Add missing fwd decl #1520

Merged

Add more forward declarations to sycl_traits.h #1497

Merged

MikeDvorskiy closed this May 7, 2024

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513

[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513

MikeDvorskiy commented Apr 18, 2024 •

edited

Loading

danhoeflinger left a comment

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

danhoeflinger Apr 18, 2024

MikeDvorskiy commented May 7, 2024

	struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};
	struct __are_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {};

[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513

[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513

Conversation

MikeDvorskiy commented Apr 18, 2024 • edited Loading

danhoeflinger left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MikeDvorskiy commented May 7, 2024

MikeDvorskiy commented Apr 18, 2024 •

edited

Loading