-
Notifications
You must be signed in to change notification settings - Fork 113
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.
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
[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types #1513
Conversation
…or transform_reduce
…or brick_copy_n and zip_iterator
0092c9b
to
3be871d
Compare
…EVICE_COPYABLE definition into the global namespace
fa945ba
to
0ed2d1b
Compare
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
|
||
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, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
} //__internal | ||
} //dpl | ||
} //oneapi |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
include/oneapi/dpl/pstl/utils.h
Outdated
@@ -48,6 +48,30 @@ namespace dpl | |||
namespace __internal | |||
{ | |||
|
|||
#if _ONEDPL_BACKEND_SYCL | |||
|
|||
template <typename... _Ts> |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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."
template <typename... _Ts> | ||
struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {}; | ||
|
||
#if __INTEL_LLVM_COMPILER && (__INTEL_LLVM_COMPILER < 20240100) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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> {};
//
include/oneapi/dpl/pstl/utils.h
Outdated
# define _ONEDPL_IS_DEVICE_COPYABLE(TYPE, ...) sycl::is_device_copyable<TYPE<__VA_ARGS__>> | ||
#endif //__INTEL_LLVM_COMPILER && (__INTEL_LLVM_COMPILER < 20240100) | ||
|
||
#else |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
#else | |
#else // _ONEDPL_BACKEND_SYCL |
include/oneapi/dpl/pstl/utils.h
Outdated
#if _ONEDPL_BACKEND_SYCL | ||
|
||
template <typename... _Ts> | ||
struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {}; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
struct __is_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {}; | |
struct __are_types_device_copyable: std::conjunction<sycl::is_device_copyable<_Ts>...> {}; |
d5b1d78
to
0539960
Compare
|
||
_ONEDPL_IS_DEVICE_COPYABLE() |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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.
@@ -375,6 +375,8 @@ struct __brick_copy_n<__hetero_tag<_BackendTag>, _ExecutionPolicy> | |||
{ | |||
__target = ::std::forward<_SourceT>(__source); | |||
} | |||
|
|||
_ONEDPL_IS_DEVICE_COPYABLE() |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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
0539960
to
f4c12b1
Compare
Close the draft because the other solution (right solution) was merged. See #1521 for details. |
[oneDPL][hetero] Re-woring sycl::is_device_copyable trait specializations for oneDPL types.
This is a prototype for 3 different oneDPL types:
transform_reduce
,__brick_copy_n
andzip_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