Refactor parallel_for serialization and fix issue caused by new XRT

[Ralender]

I will add a test for this.

[fxzjshm]

Sorry if I disturb you, but I have some problem using this branch.
The current implementation requires the kernel to be device copyable, which shouldn't be required (I think, at least non-device-copyable kernels compiles on other SYCL implementations.)

Consider this situation: a kernel K uses a parameter P that isn't device copyable because it isn't "TriviallyCopyable" (in my case, P's type is boost::iterators::counting_iterator<T>), but I think it should work on device, so I explicitly added template <typename T> struct sycl::is_device_copyable<boost::iterators::counting_iterator<T>> : std::true_type {};, and it really works. However, this parameter P makes the kernel K itself non-"TriviallyCopyable" and non-device-copyable, which is required in this implemention (maybe by detail::serialize_parallel_for() here and here as the kernel K becomes detail::serialize_parallel_for()'s parameter). In my humble opinion, It's hard to deal with this using explicit sycl::is_device_copyable<...> as the kernel is often a lambda.

Could you help me how to deal with this problem?

[Ralender]

I pushed a commit that should address you issue. becasue we should respect user provided specializations of sycl::is_device_copyable. I also noticed that it doesn't seem like there is any checking that the lambda only contains device copyable type. for parallel_for.

but it is rarely safe to bitwise copy to the device a type that are not TriviallyCopyable. especially an iterator which is likely to contain a pointer. I didn't look into boost::iterators::counting_iterator to check if it would be safe.

[fxzjshm]

Yes, in most cases you are right. But in some edge cases it just don't work 😢 (e.g. std::pair-like user-defined templated class with custom copy assignment operator, even if that operator= does the same thing as = default when T is TriviallyCopyable)

Anyway, thanks for this!