Removal ExecutionPolicy from submitters and etc. stuff

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h

@@ -25,7 +25,7 @@ namespace oneapi::dpl::experimental::kt::gpu::esimd
 // TODO: make sure to provide sufficient diagnostic if input does not allow either reading or writing
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng>
 std::enable_if_t<!oneapi::dpl::__internal::__is_iterator_type_v<_KeysRng>, sycl::event>
-radix_sort(sycl::queue __q, _KeysRng&& __keys_rng, _KernelParam __param = {})
+radix_sort(sycl::queue& __q, _KeysRng&& __keys_rng, _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
 
@@ -38,7 +38,7 @@ radix_sort(sycl::queue __q, _KeysRng&& __keys_rng, _KernelParam __param = {})
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator>
 std::enable_if_t<oneapi::dpl::__internal::__is_iterator_type_v<_KeysIterator>, sycl::event>
-radix_sort(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __keys_last, _KernelParam __param = {})
+radix_sort(sycl::queue& __q, _KeysIterator __keys_first, _KeysIterator __keys_last, _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
 
@@ -54,7 +54,7 @@ radix_sort(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __keys_las
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng,
           typename _ValsRng>
 std::enable_if_t<!oneapi::dpl::__internal::__is_iterator_type_v<_KeysRng>, sycl::event>
-radix_sort_by_key(sycl::queue __q, _KeysRng&& __keys_rng, _ValsRng&& __vals_rng, _KernelParam __param = {})
+radix_sort_by_key(sycl::queue& __q, _KeysRng&& __keys_rng, _ValsRng&& __vals_rng, _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
 
@@ -69,7 +69,7 @@ radix_sort_by_key(sycl::queue __q, _KeysRng&& __keys_rng, _ValsRng&& __vals_rng,
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator,
           typename _ValsIterator>
 std::enable_if_t<oneapi::dpl::__internal::__is_iterator_type_v<_KeysIterator>, sycl::event>
-radix_sort_by_key(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __keys_last, _ValsIterator __vals_first,
+radix_sort_by_key(sycl::queue& __q, _KeysIterator __keys_first, _KeysIterator __keys_last, _ValsIterator __vals_first,
                   _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
@@ -89,7 +89,7 @@ radix_sort_by_key(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __k
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng1,
           typename _KeysRng2>
 std::enable_if_t<!oneapi::dpl::__internal::__is_iterator_type_v<_KeysRng1>, sycl::event>
-radix_sort(sycl::queue __q, _KeysRng1&& __keys_rng, _KeysRng2&& __keys_rng_out, _KernelParam __param = {})
+radix_sort(sycl::queue& __q, _KeysRng1&& __keys_rng, _KeysRng2&& __keys_rng_out, _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
     if (__keys_rng.size() == 0)
@@ -104,7 +104,7 @@ radix_sort(sycl::queue __q, _KeysRng1&& __keys_rng, _KeysRng2&& __keys_rng_out,
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator1,
           typename _KeysIterator2>
 std::enable_if_t<oneapi::dpl::__internal::__is_iterator_type_v<_KeysIterator1>, sycl::event>
-radix_sort(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_last, _KeysIterator2 __keys_out_first,
+radix_sort(sycl::queue& __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_last, _KeysIterator2 __keys_out_first,
            _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
@@ -126,7 +126,7 @@ radix_sort(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_l
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng1,
           typename _ValsRng1, typename _KeysRng2, typename _ValsRng2>
 std::enable_if_t<!oneapi::dpl::__internal::__is_iterator_type_v<_KeysRng1>, sycl::event>
-radix_sort_by_key(sycl::queue __q, _KeysRng1&& __keys_rng, _ValsRng1&& __vals_rng, _KeysRng2&& __keys_out_rng,
+radix_sort_by_key(sycl::queue& __q, _KeysRng1&& __keys_rng, _ValsRng1&& __vals_rng, _KeysRng2&& __keys_out_rng,
                   _ValsRng2&& __vals_out_rng, _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
@@ -144,8 +144,9 @@ radix_sort_by_key(sycl::queue __q, _KeysRng1&& __keys_rng, _ValsRng1&& __vals_rn
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator1,
           typename _ValsIterator1, typename _KeysIterator2, typename _ValsIterator2>
 std::enable_if_t<oneapi::dpl::__internal::__is_iterator_type_v<_KeysIterator1>, sycl::event>
-radix_sort_by_key(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_last, _ValsIterator1 __vals_first,
-                  _KeysIterator2 __keys_out_first, _ValsIterator2 __vals_out_first, _KernelParam __param = {})
+radix_sort_by_key(sycl::queue& __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_last,
+                  _ValsIterator1 __vals_first, _KeysIterator2 __keys_out_first, _ValsIterator2 __vals_out_first,
+                  _KernelParam __param = {})
 {
     __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
 

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

@@ -51,7 +51,7 @@ class __esimd_radix_sort_onesweep_copyback_by_key;
 template <typename _KernelName, bool __is_ascending, ::std::uint8_t __radix_bits, ::std::uint16_t __data_per_work_item,
           ::std::uint16_t __work_group_size, typename _RngPack>
 sycl::event
-__one_wg(sycl::queue __q, _RngPack&& __pack, ::std::size_t __n)
+__one_wg(sycl::queue& __q, _RngPack&& __pack, std::size_t __n)
 {
     using _KeyT = typename ::std::decay_t<_RngPack>::_KeyT;
     using _EsimRadixSortKernel =
@@ -64,7 +64,7 @@ __one_wg(sycl::queue __q, _RngPack&& __pack, ::std::size_t __n)
 template <typename _KernelName, bool __is_ascending, ::std::uint8_t __radix_bits, ::std::uint16_t __data_per_work_item,
           ::std::uint16_t __work_group_size, typename _RngPack1, typename _RngPack2>
 sycl::event
-__one_wg(sycl::queue __q, _RngPack1&& __pack_in, _RngPack2&& __pack_out, ::std::size_t __n)
+__one_wg(sycl::queue& __q, _RngPack1&& __pack_in, _RngPack2&& __pack_out, std::size_t __n)
 {
     using _KeyT = typename ::std::decay_t<_RngPack1>::_KeyT;
     using _EsimRadixSortKernel =

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_submitters.h

@@ -41,7 +41,7 @@ struct __radix_sort_one_wg_submitter<__is_ascending, __radix_bits, __data_per_wo
 {
     template <typename _RngPack1, typename _RngPack2>
     sycl::event
-    operator()(sycl::queue __q, _RngPack1&& __pack_in, _RngPack2&& __pack_out, ::std::size_t __n) const
+    operator()(sycl::queue& __q, _RngPack1&& __pack_in, _RngPack2&& __pack_out, std::size_t __n) const
     {
         sycl::nd_range<1> __nd_range{__work_group_size, __work_group_size};
         return __q.submit([&](sycl::handler& __cgh) {

include/oneapi/dpl/experimental/kt/single_pass_scan.h

@@ -145,7 +145,7 @@ struct __lookback_init_submitter<_FlagType, _Type, _BinaryOp,
 {
     template <typename _StatusFlags, typename _PartialValues>
     sycl::event
-    operator()(sycl::queue __q, _StatusFlags&& __status_flags, _PartialValues&& __partial_values,
+    operator()(sycl::queue& __q, _StatusFlags&& __status_flags, _PartialValues&& __partial_values,
                std::size_t __status_flags_size, std::uint16_t __status_flag_padding) const
     {
         return __q.submit([&](sycl::handler& __hdl) {
@@ -276,8 +276,8 @@ struct __lookback_submitter<__data_per_workitem, __workgroup_size, _Type, _FlagT
 
     template <typename _InRng, typename _OutRng, typename _BinaryOp, typename _StatusFlags, typename _StatusValues>
     sycl::event
-    operator()(sycl::queue __q, sycl::event __prev_event, _InRng&& __in_rng, _OutRng&& __out_rng, _BinaryOp __binary_op,
-               std::size_t __n, _StatusFlags&& __status_flags, std::size_t __status_flags_size,
+    operator()(sycl::queue& __q, sycl::event __prev_event, _InRng&& __in_rng, _OutRng&& __out_rng,
+               _BinaryOp __binary_op, std::size_t __n, _StatusFlags&& __status_flags, std::size_t __status_flags_size,
                _StatusValues&& __status_vals_full, _StatusValues&& __status_vals_partial,
                std::size_t __current_num_items) const
     {
@@ -304,7 +304,7 @@ struct __lookback_submitter<__data_per_workitem, __workgroup_size, _Type, _FlagT
 
 template <bool _Inclusive, typename _InRange, typename _OutRange, typename _BinaryOp, typename _KernelParam>
 sycl::event
-__single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_rng, _BinaryOp __binary_op, _KernelParam)
+__single_pass_scan(sycl::queue& __q, _InRange&& __in_rng, _OutRange&& __out_rng, _BinaryOp __binary_op, _KernelParam)
 {
     using _Type = oneapi::dpl::__internal::__value_t<_InRange>;
     using _FlagType = __scan_status_flag<_Type>;
@@ -326,19 +326,19 @@ __single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_r
                   "Only binary operators with known identity values are supported");
 
     assert("This device does not support 64-bit atomics" &&
-           (sizeof(_Type) < 8 || __queue.get_device().has(sycl::aspect::atomic64)));
+           (sizeof(_Type) < 8 || __q.get_device().has(sycl::aspect::atomic64)));
 
     // Next power of 2 greater than or equal to __n
     auto __n_uniform = ::oneapi::dpl::__internal::__dpl_bit_ceil(__n);
 
     // Perform a single-work group scan if the input is small
-    if (oneapi::dpl::__par_backend_hetero::__group_scan_fits_in_slm<_Type>(__queue, __n, __n_uniform, /*limit=*/16384))
+    if (oneapi::dpl::__par_backend_hetero::__group_scan_fits_in_slm<_Type>(__q, __n, __n_uniform, /*limit=*/16384))
     {
-        return oneapi::dpl::__par_backend_hetero::__parallel_transform_scan_single_group(
-            oneapi::dpl::__internal::__device_backend_tag{},
-            oneapi::dpl::execution::__dpl::make_device_policy<typename _KernelParam::kernel_name>(__queue),
-            std::forward<_InRange>(__in_rng), std::forward<_OutRange>(__out_rng), __n,
-            oneapi::dpl::__internal::__no_op{}, unseq_backend::__no_init_value<_Type>{}, __binary_op, std::true_type{});
+        return oneapi::dpl::__par_backend_hetero::__parallel_transform_scan_single_group<
+            typename _KernelParam::kernel_name>(oneapi::dpl::__internal::__device_backend_tag{}, __q,
+                                                std::forward<_InRange>(__in_rng), std::forward<_OutRange>(__out_rng),
+                                                __n, oneapi::dpl::__internal::__no_op{},
+                                                unseq_backend::__no_init_value<_Type>{}, __binary_op, std::true_type{});
     }
 
     constexpr std::size_t __workgroup_size = _KernelParam::workgroup_size;
@@ -358,7 +358,7 @@ __single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_r
     std::size_t __mem_bytes =
         __status_flags_bytes + __status_vals_full_offset_bytes + __status_vals_partial_offset_bytes + __mem_align_pad;
 
-    std::byte* __device_mem = reinterpret_cast<std::byte*>(sycl::malloc_device(__mem_bytes, __queue));
+    std::byte* __device_mem = reinterpret_cast<std::byte*>(sycl::malloc_device(__mem_bytes, __q));
     if (!__device_mem)
         throw std::bad_alloc();
 
@@ -372,14 +372,14 @@ __single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_r
         reinterpret_cast<_Type*>(__status_vals_full + __status_vals_full_offset_bytes / sizeof(_Type));
 
     auto __fill_event = __lookback_init_submitter<_FlagType, _Type, _BinaryOp, _LookbackInitKernel>{}(
-        __queue, __status_flags, __status_vals_partial, __status_flags_size, __status_flag_padding);
+        __q, __status_flags, __status_vals_partial, __status_flags_size, __status_flag_padding);
 
     std::size_t __current_num_wgs = oneapi::dpl::__internal::__dpl_ceiling_div(__n, __elems_in_tile);
     std::size_t __current_num_items = __current_num_wgs * __workgroup_size;
 
     auto __prev_event =
         __lookback_submitter<__data_per_workitem, __workgroup_size, _Type, _FlagType, _LookbackKernel>{}(
-            __queue, __fill_event, __in_rng, __out_rng, __binary_op, __n, __status_flags, __status_flags_size,
+            __q, __fill_event, __in_rng, __out_rng, __binary_op, __n, __status_flags, __status_flags_size,
             __status_vals_full, __status_vals_partial, __current_num_items);
 
     // TODO: Currently, the following portion of code makes this entire function synchronous.
@@ -388,15 +388,15 @@ __single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_r
     // we should replace this code with the asynchronous version below.
     if (0)
     {
-        return __queue.submit([=](sycl::handler& __hdl) {
+        return __q.submit([=](sycl::handler& __hdl) {
             __hdl.depends_on(__prev_event);
-            __hdl.host_task([=]() { sycl::free(__device_mem, __queue); });
+            __hdl.host_task([=]() { sycl::free(__device_mem, __q); });
         });
     }
     else
     {
         __prev_event.wait();
-        sycl::free(__device_mem, __queue);
+        sycl::free(__device_mem, __q);
         return __prev_event;
     }
 }
@@ -405,18 +405,18 @@ __single_pass_scan(sycl::queue __queue, _InRange&& __in_rng, _OutRange&& __out_r
 
 template <typename _InRng, typename _OutRng, typename _BinaryOp, typename _KernelParam>
 sycl::event
-inclusive_scan(sycl::queue __queue, _InRng&& __in_rng, _OutRng&& __out_rng, _BinaryOp __binary_op,
+inclusive_scan(sycl::queue& __q, _InRng&& __in_rng, _OutRng&& __out_rng, _BinaryOp __binary_op,
                _KernelParam __param = {})
 {
     auto __in_view = oneapi::dpl::__ranges::views::all(std::forward<_InRng>(__in_rng));
     auto __out_view = oneapi::dpl::__ranges::views::all(std::forward<_OutRng>(__out_rng));
 
-    return __impl::__single_pass_scan<true>(__queue, std::move(__in_view), std::move(__out_view), __binary_op, __param);
+    return __impl::__single_pass_scan<true>(__q, std::move(__in_view), std::move(__out_view), __binary_op, __param);
 }
 
 template <typename _InIterator, typename _OutIterator, typename _BinaryOp, typename _KernelParam>
 sycl::event
-inclusive_scan(sycl::queue __queue, _InIterator __in_begin, _InIterator __in_end, _OutIterator __out_begin,
+inclusive_scan(sycl::queue& __q, _InIterator __in_begin, _InIterator __in_end, _OutIterator __out_begin,
                _BinaryOp __binary_op, _KernelParam __param = {})
 {
     auto __n = __in_end - __in_begin;
@@ -426,7 +426,7 @@ inclusive_scan(sycl::queue __queue, _InIterator __in_begin, _InIterator __in_end
     auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _OutIterator>();
     auto __buf2 = __keep2(__out_begin, __out_begin + __n);
 
-    return __impl::__single_pass_scan<true>(__queue, __buf1.all_view(), __buf2.all_view(), __binary_op, __param);
+    return __impl::__single_pass_scan<true>(__q, __buf1.all_view(), __buf2.all_view(), __binary_op, __param);
 }
 
 } // namespace gpu

include/oneapi/dpl/internal/async_impl/async_impl_hetero.h

@@ -140,9 +140,13 @@ __pattern_transform_reduce_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& _
         oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _RandomAccessIterator2>();
     auto __buf2 = __keep2(__first2, __first2 + __n);
 
-    return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_RepackedTp,
-                                                                          ::std::true_type /*is_commutative*/>(
-        _BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __binary_op1, _Functor{__binary_op2},
+    using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
+
+    sycl::queue __q_local = __exec.queue();
+
+    return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_CustomName, _RepackedTp,
+                                                                          std::true_type /*is_commutative*/>(
+        _BackendTag{}, __q_local, __binary_op1, _Functor{__binary_op2},
         unseq_backend::__init_value<_RepackedTp>{__init}, // initial value
         __buf1.all_view(), __buf2.all_view());
 }
@@ -166,9 +170,13 @@ __pattern_transform_reduce_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& _
     auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _ForwardIterator>();
     auto __buf = __keep(__first, __last);
 
-    return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_RepackedTp,
-                                                                          ::std::true_type /*is_commutative*/>(
-        _BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __binary_op, _Functor{__unary_op},
+    using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
+
+    sycl::queue __q_local = __exec.queue();
+
+    return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_CustomName, _RepackedTp,
+                                                                          std::true_type /*is_commutative*/>(
+        _BackendTag{}, __q_local, __binary_op, _Functor{__unary_op},
         unseq_backend::__init_value<_RepackedTp>{__init}, // initial value
         __buf.all_view());
 }
@@ -204,9 +212,13 @@ __pattern_transform_scan_base_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&
     auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _Iterator2>();
     auto __buf2 = __keep2(__result, __result + __n);
 
-    auto __res = oneapi::dpl::__par_backend_hetero::__parallel_transform_scan(
-        _BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec), __buf1.all_view(), __buf2.all_view(), __n, __unary_op,
-        __init, __binary_op, _Inclusive{});
+    sycl::queue __q_local = __exec.queue();
+
+    using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
+
+    auto __res = oneapi::dpl::__par_backend_hetero::__parallel_transform_scan<_CustomName>(
+        _BackendTag{}, __q_local, __buf1.all_view(), __buf2.all_view(), __n, __unary_op, __init, __binary_op,
+        _Inclusive{});
     return __res.__make_future(__result + __n);
 }
 

include/oneapi/dpl/internal/async_impl/glue_async_impl.h

@@ -102,8 +102,12 @@ sort_async(_ExecutionPolicy&& __exec, _Iterator __first, _Iterator __last, _Comp
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
     using __backend_tag = typename decltype(__dispatch_tag)::__backend_tag;
 
-    return __par_backend_hetero::__parallel_stable_sort(__backend_tag{}, ::std::forward<_ExecutionPolicy>(__exec),
-                                                        __buf.all_view(), __comp, oneapi::dpl::identity{});
+    sycl::queue __q_local = __exec.queue();
+
+    using _CustomName = oneapi::dpl::__internal::__policy_kernel_name<_ExecutionPolicy>;
+
+    return __par_backend_hetero::__parallel_stable_sort<_CustomName>(__backend_tag{}, __q_local, __buf.all_view(),
+                                                                     __comp, oneapi::dpl::identity{});
 }
 
 template <class _ExecutionPolicy, class _RandomAccessIterator, class... _Events,