Implement a SYCL Onesweep Radix Sort KT

include/oneapi/dpl/experimental/kernel_templates

@@ -18,6 +18,7 @@
 #    include "kt/esimd_radix_sort.h"
 #endif
 
+#include "kt/sycl_radix_sort.h"
 #include "kt/single_pass_scan.h"
 
 #endif // _ONEDPL_KERNEL_TEMPLATES

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

@@ -16,7 +16,8 @@
 #include "../../pstl/hetero/dpcpp/utils_ranges_sycl.h"
 
 #include "internal/esimd_radix_sort_utils.h"
-#include "internal/esimd_radix_sort_dispatchers.h"
+#include "internal/radix_sort_utils.h"
+#include "internal/radix_sort_dispatchers.h"
 #include "../../pstl/utils.h"
 
 namespace oneapi::dpl::experimental::kt::gpu::esimd
@@ -33,7 +34,8 @@ radix_sort(sycl::queue __q, _KeysRng&& __keys_rng, _KernelParam __param = {})
         return {};
 
     auto __pack = __impl::__rng_pack{oneapi::dpl::__ranges::views::all(::std::forward<_KeysRng>(__keys_rng))};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, __pack, __pack, __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator>
@@ -48,7 +50,8 @@ radix_sort(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __keys_las
     auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
     auto __keys_rng = __keys_keep(__keys_first, __keys_last).all_view();
     auto __pack = __impl::__rng_pack{::std::move(__keys_rng)};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, __pack, __pack, __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng,
@@ -63,7 +66,8 @@ radix_sort_by_key(sycl::queue __q, _KeysRng&& __keys_rng, _ValsRng&& __vals_rng,
 
     auto __pack = __impl::__rng_pack{oneapi::dpl::__ranges::views::all(::std::forward<_KeysRng>(__keys_rng)),
                                      oneapi::dpl::__ranges::views::all(::std::forward<_ValsRng>(__vals_rng))};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, __pack, __pack, __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator,
@@ -83,7 +87,8 @@ radix_sort_by_key(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __k
     auto __vals_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
     auto __vals_rng = __vals_keep(__vals_first, __vals_first + (__keys_last - __keys_first)).all_view();
     auto __pack = __impl::__rng_pack{::std::move(__keys_rng), ::std::move(__vals_rng)};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, __pack, __pack, __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng1,
@@ -97,8 +102,9 @@ radix_sort(sycl::queue __q, _KeysRng1&& __keys_rng, _KeysRng2&& __keys_rng_out,
 
     auto __pack = __impl::__rng_pack{oneapi::dpl::__ranges::views::all(::std::forward<_KeysRng1>(__keys_rng))};
     auto __pack_out = __impl::__rng_pack{oneapi::dpl::__ranges::views::all(::std::forward<_KeysRng2>(__keys_rng_out))};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),
-                                                                                    ::std::move(__pack_out), __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, std::move(__pack), std::move(__pack_out),
+        __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator1,
@@ -119,8 +125,9 @@ radix_sort(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_l
     auto __keys_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
     auto __keys_out_rng = __keys_out_keep(__keys_out_first, __keys_out_first + __n).all_view();
     auto __pack_out = __impl::__rng_pack{::std::move(__keys_out_rng)};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),
-                                                                                    ::std::move(__pack_out), __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, std::move(__pack), std::move(__pack_out),
+        __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysRng1,
@@ -137,8 +144,9 @@ radix_sort_by_key(sycl::queue __q, _KeysRng1&& __keys_rng, _ValsRng1&& __vals_rn
                                      oneapi::dpl::__ranges::views::all(::std::forward<_ValsRng1>(__vals_rng))};
     auto __pack_out = __impl::__rng_pack{oneapi::dpl::__ranges::views::all(::std::forward<_KeysRng2>(__keys_out_rng)),
                                          oneapi::dpl::__ranges::views::all(::std::forward<_ValsRng2>(__vals_out_rng))};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),
-                                                                                    ::std::move(__pack_out), __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, std::move(__pack), std::move(__pack_out),
+        __param);
 }
 
 template <bool __is_ascending = true, ::std::uint8_t __radix_bits = 8, typename _KernelParam, typename _KeysIterator1,
@@ -166,8 +174,9 @@ radix_sort_by_key(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 _
     auto __vals_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
     auto __vals_out_rng = __vals_out_keep(__vals_out_first, __vals_out_first + __n).all_view();
     auto __pack_out = __impl::__rng_pack{::std::move(__keys_out_rng), ::std::move(__vals_out_rng)};
-    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),
-                                                                                    ::std::move(__pack_out), __param);
+    return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(
+        oneapi::dpl::experimental::kt::gpu::__impl::__esimd_tag{}, __q, std::move(__pack), std::move(__pack_out),
+        __param);
 }
 
 } // namespace oneapi::dpl::experimental::kt::gpu::esimd

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

@@ -23,13 +23,13 @@
 
 #define _ONEDPL_ESIMD_LSC_FENCE_PRESENT (_ONEDPL_LIBSYCL_VERSION >= 70200)
 
-namespace oneapi::dpl::experimental::kt::gpu::esimd::__impl
+namespace oneapi::dpl::experimental::kt::gpu::__impl
 {
 
 // TODO: rename to show the meaning clearly: default vectorization factor
 constexpr int __data_per_step = 16;
 
-} // namespace oneapi::dpl::experimental::kt::gpu::esimd::__impl
+} // namespace oneapi::dpl::experimental::kt::gpu::__impl
 
 // This namespace mostly consists of abstractions on the top of regular ESIMD functions.
 // The purpose is: