[SYCL][InvokeSIMD] Add error for invalid uniform arguments #8916
Conversation
| @@ -420,6 +431,7 @@ __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg, | |||
| // what the subgroup size is and arguments don't need widening and return | |||
| // value does not need shrinking by this library or SPMD compiler, so 0 | |||
| // is fine in this case. | |||
| detail::verify_valid_args<T...>(); | |||
| detail::verify_no_ref<Callable>(); | |||
There was a problem hiding this comment.
The name of the function seems a bit too general. Now it checks only that all arguments are trivially_copiable.
If keep the name, then it would be reasonable to move checks from verify_no_ref to that function as well.
There was a problem hiding this comment.
good point, im gonna pick the second option so it is more extensible in the future, thanks
| template <typename T> struct is_non_trivially_copyable_uniform { | ||
| static constexpr bool value = | ||
| is_uniform_type<T>::value && | ||
| !std::is_trivially_copyable_v<typename unwrap_uniform<T>::type>; |
There was a problem hiding this comment.
@rolandschulz , @Pennycook - is the trivially-copyable requirement applied to parameters ONLY ON DEVICE?
Let's say, sycl::local_accessor has different representation on HOST and DEVICE. It uses std::shared_ptr on HOST and simpler types on DEVICE. Such object is trivially_copyable on DEVICE
So, should this check be applied like this?
#ifdef __SYCL_DEVICE_ONLY__
check that invoke_simd arg is trivially copyable
#endif
|
@v-klochkov Should I change it so this check only happens for device code, and then we can merge? Yes, please do the trivially-copyable check under |
The spec requires uniform arguments to contain a trivially copyable type. Add an error if it isn't. Signed-off-by: Sarnie, Nick <[email protected]>
Signed-off-by: Sarnie, Nick <[email protected]>
Signed-off-by: Sarnie, Nick <[email protected]>
Signed-off-by: Sarnie, Nick <[email protected]>
|
@v-klochkov ok this is ready for you to take another look, thanks! |
|
The spec requires uniform arguments to contain a trivially copyable type. Add an error if it isn't.