| 297 | |
| 298 | template <class ForwardIt, class T> |
| 299 | sycl::event fill(sycl::queue &q, ForwardIt first, ForwardIt last, |
| 300 | const T &value, const std::vector<sycl::event> &deps = {}) { |
| 301 | auto size = std::distance(first, last); |
| 302 | if(size == 0) |
| 303 | return sycl::event{}; |
| 304 | |
| 305 | using value_type = typename std::iterator_traits<ForwardIt>::value_type; |
| 306 | |
| 307 | auto invoke_kernel = [&]() -> sycl::event{ |
| 308 | return q.parallel_for(sycl::range{size}, deps, |
| 309 | [=](sycl::id<1> id) { |
| 310 | auto it = first; |
| 311 | std::advance(it, id[0]); |
| 312 | *it = value; |
| 313 | }); |
| 314 | }; |
| 315 | |
| 316 | if constexpr (std::is_trivial_v<value_type> && |
| 317 | std::is_same_v<value_type, T> && |
| 318 | util::is_contiguous<ForwardIt>()) { |
| 319 | unsigned char equal_byte; |
| 320 | if (detail::all_bytes_equal(value, equal_byte) && |
| 321 | detail::should_use_memset(q.get_device())) { |
| 322 | return q.memset(&(*first), static_cast<int>(equal_byte), |
| 323 | size * sizeof(T), deps); |
| 324 | } else { |
| 325 | return invoke_kernel(); |
| 326 | } |
| 327 | } else { |
| 328 | return invoke_kernel(); |
| 329 | } |
| 330 | } |
| 331 | |
| 332 | template<class ForwardIt, class Size, class T > |
| 333 | sycl::event fill_n(sycl::queue& q, |