Skip to content

Commit c4d223e

Browse files
danhoeflingerCopilot
danhoeflinger
and
Copilot
authored
Refactor __get_sycl_range to align with sycl (#2519)
--------- Signed-off-by: Dan Hoeflinger <dan.hoeflinger@intel.com> Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
1 parent 21474c0 commit c4d223e

20 files changed

+583
-318
lines changed

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

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -45,7 +45,7 @@ radix_sort(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __keys_las
4545
if (__keys_last - __keys_first < 2)
4646
return {};
4747

48-
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator>();
48+
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
4949
auto __keys_rng = __keys_keep(__keys_first, __keys_last).all_view();
5050
auto __pack = __impl::__rng_pack{::std::move(__keys_rng)};
5151
return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
@@ -77,10 +77,10 @@ radix_sort_by_key(sycl::queue __q, _KeysIterator __keys_first, _KeysIterator __k
7777
if (__keys_last - __keys_first < 2)
7878
return {};
7979

80-
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator>();
80+
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
8181
auto __keys_rng = __keys_keep(__keys_first, __keys_last).all_view();
8282

83-
auto __vals_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _ValsIterator>();
83+
auto __vals_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
8484
auto __vals_rng = __vals_keep(__vals_first, __vals_first + (__keys_last - __keys_first)).all_view();
8585
auto __pack = __impl::__rng_pack{::std::move(__keys_rng), ::std::move(__vals_rng)};
8686
return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/true>(__q, __pack, __pack, __param);
@@ -113,10 +113,10 @@ radix_sort(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 __keys_l
113113
if (__n == 0)
114114
return {};
115115

116-
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator1>();
116+
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
117117
auto __keys_rng = __keys_keep(__keys_first, __keys_last).all_view();
118118
auto __pack = __impl::__rng_pack{::std::move(__keys_rng)};
119-
auto __keys_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator2>();
119+
auto __keys_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
120120
auto __keys_out_rng = __keys_out_keep(__keys_out_first, __keys_out_first + __n).all_view();
121121
auto __pack_out = __impl::__rng_pack{::std::move(__keys_out_rng)};
122122
return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),
@@ -153,17 +153,17 @@ radix_sort_by_key(sycl::queue __q, _KeysIterator1 __keys_first, _KeysIterator1 _
153153
if (__n == 0)
154154
return {};
155155

156-
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator1>();
156+
auto __keys_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
157157
auto __keys_rng = __keys_keep(__keys_first, __keys_last).all_view();
158158

159-
auto __vals_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _ValsIterator1>();
159+
auto __vals_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
160160
auto __vals_rng = __vals_keep(__vals_first, __vals_first + __n).all_view();
161161
auto __pack = __impl::__rng_pack{::std::move(__keys_rng), ::std::move(__vals_rng)};
162162

163-
auto __keys_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeysIterator2>();
163+
auto __keys_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
164164
auto __keys_out_rng = __keys_out_keep(__keys_out_first, __keys_out_first + __n).all_view();
165165

166-
auto __vals_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _ValsIterator2>();
166+
auto __vals_out_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
167167
auto __vals_out_rng = __vals_out_keep(__vals_out_first, __vals_out_first + __n).all_view();
168168
auto __pack_out = __impl::__rng_pack{::std::move(__keys_out_rng), ::std::move(__vals_out_rng)};
169169
return __impl::__radix_sort<__is_ascending, __radix_bits, /*__in_place=*/false>(__q, ::std::move(__pack),

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -328,12 +328,12 @@ __onesweep(sycl::queue __q, _RngPack1&& __pack, _RngPack2&& __pack_out, ::std::s
328328
__mem_holder.__allocate();
329329

330330
auto __get_tmp_pack = [&]() {
331-
auto __keys_tmp_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _KeyT*>();
331+
auto __keys_tmp_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
332332
auto __keys_tmp_rng = __keys_tmp_keep(__mem_holder.__keys_ptr(), __mem_holder.__keys_ptr() + __n).all_view();
333333

334334
if constexpr (__has_values)
335335
{
336-
auto __vals_tmp_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write, _ValT*>();
336+
auto __vals_tmp_keep = oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read_write>();
337337
auto __vals_tmp_rng =
338338
__vals_tmp_keep(__mem_holder.__vals_ptr(), __mem_holder.__vals_ptr() + __n).all_view();
339339
return __rng_pack(::std::move(__keys_tmp_rng), ::std::move(__vals_tmp_rng));

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -326,9 +326,10 @@ inclusive_scan(sycl::queue __queue, _InIterator __in_begin, _InIterator __in_end
326326
{
327327
auto __n = __in_end - __in_begin;
328328

329-
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _InIterator>();
329+
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
330330
auto __buf1 = __keep1(__in_begin, __in_end);
331-
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _OutIterator>();
331+
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write,
332+
/*_IsNoInitRequested=*/true>();
332333
auto __buf2 = __keep2(__out_begin, __out_begin + __n);
333334

334335
return __impl::__single_pass_scan<true>(__queue, __buf1.all_view(), __buf2.all_view(), __binary_op, __param);

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

Lines changed: 21 additions & 26 deletions
Original file line numberDiff line numberDiff line change
@@ -30,16 +30,16 @@ namespace dpl
3030
namespace __internal
3131
{
3232

33-
template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
33+
template <__par_backend_hetero::access_mode __acc_mode, bool _IsNoInitRequested, typename _BackendTag,
34+
typename _ExecutionPolicy, typename _ForwardIterator, typename _Function>
3435
auto
3536
__pattern_walk1_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator __first,
3637
_ForwardIterator __last, _Function __f)
3738
{
3839
auto __n = __last - __first;
3940
assert(__n > 0);
4041

41-
auto __keep =
42-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read_write, _ForwardIterator>();
42+
auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode, _IsNoInitRequested>();
4343
auto __buf = __keep(__first, __last);
4444

4545
auto __future_obj = oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -48,21 +48,19 @@ __pattern_walk1_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _For
4848
return __future_obj;
4949
}
5050

51-
template <__par_backend_hetero::access_mode __acc_mode1 = __par_backend_hetero::access_mode::read,
52-
__par_backend_hetero::access_mode __acc_mode2 = __par_backend_hetero::access_mode::write,
53-
typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator1, typename _ForwardIterator2,
54-
typename _Function>
51+
template <__par_backend_hetero::access_mode __out_acc_mode, bool _IsOutNoInitRequested, typename _BackendTag,
52+
typename _ExecutionPolicy, typename _ForwardIterator1, typename _ForwardIterator2, typename _Function>
5553
auto
5654
__pattern_walk2_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator1 __first1,
5755
_ForwardIterator1 __last1, _ForwardIterator2 __first2, _Function __f)
5856
{
5957
auto __n = __last1 - __first1;
6058
assert(__n > 0);
6159

62-
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode1, _ForwardIterator1>();
60+
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<sycl::access::mode::read>();
6361
auto __buf1 = __keep1(__first1, __last1);
6462

65-
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__acc_mode2, _ForwardIterator2>();
63+
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__out_acc_mode, _IsOutNoInitRequested>();
6664
auto __buf2 = __keep2(__first2, __first2 + __n);
6765

6866
auto __future = oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -73,23 +71,21 @@ __pattern_walk2_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _For
7371
return __future.__make_future(__first2 + __n);
7472
}
7573

76-
template <typename _BackendTag, typename _ExecutionPolicy, typename _ForwardIterator1, typename _ForwardIterator2,
77-
typename _ForwardIterator3, typename _Function>
74+
template <__par_backend_hetero::access_mode __output_acc_mode, bool _IsOutNoInitRequested, typename _BackendTag,
75+
typename _ExecutionPolicy, typename _ForwardIterator1, typename _ForwardIterator2, typename _ForwardIterator3,
76+
typename _Function>
7877
auto
7978
__pattern_walk3_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& __exec, _ForwardIterator1 __first1,
8079
_ForwardIterator1 __last1, _ForwardIterator2 __first2, _ForwardIterator3 __first3, _Function __f)
8180
{
8281
auto __n = __last1 - __first1;
8382
assert(__n > 0);
8483

85-
auto __keep1 =
86-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _ForwardIterator1>();
84+
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
8785
auto __buf1 = __keep1(__first1, __last1);
88-
auto __keep2 =
89-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _ForwardIterator2>();
86+
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
9087
auto __buf2 = __keep2(__first2, __first2 + __n);
91-
auto __keep3 =
92-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _ForwardIterator3>();
88+
auto __keep3 = oneapi::dpl::__ranges::__get_sycl_range<__output_acc_mode, _IsOutNoInitRequested>();
9389
auto __buf3 = __keep3(__first3, __first3 + __n);
9490

9591
auto __future = oneapi::dpl::__par_backend_hetero::__parallel_for(
@@ -106,7 +102,7 @@ auto
106102
__pattern_walk2_brick_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator1 __first1,
107103
_ForwardIterator1 __last1, _ForwardIterator2 __first2, _Brick __brick)
108104
{
109-
return __pattern_walk2_async(
105+
return __pattern_walk2_async<__par_backend_hetero::access_mode::write, /*_IsNoInitRequested=*/true>(
110106
__tag,
111107
__par_backend_hetero::make_wrapped_policy<__walk2_brick_wrapper>(::std::forward<_ExecutionPolicy>(__exec)),
112108
__first1, __last1, __first2, __brick);
@@ -129,11 +125,9 @@ __pattern_transform_reduce_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& _
129125
using _RepackedTp = __par_backend_hetero::__repacked_tuple_t<_Tp>;
130126

131127
auto __n = __last1 - __first1;
132-
auto __keep1 =
133-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _RandomAccessIterator1>();
128+
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
134129
auto __buf1 = __keep1(__first1, __last1);
135-
auto __keep2 =
136-
oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _RandomAccessIterator2>();
130+
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
137131
auto __buf2 = __keep2(__first2, __first2 + __n);
138132

139133
return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_RepackedTp,
@@ -159,7 +153,7 @@ __pattern_transform_reduce_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&& _
159153
using _Functor = unseq_backend::walk_n<_UnaryOperation>;
160154
using _RepackedTp = __par_backend_hetero::__repacked_tuple_t<_Tp>;
161155

162-
auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _ForwardIterator>();
156+
auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
163157
auto __buf = __keep(__first, __last);
164158

165159
return oneapi::dpl::__par_backend_hetero::__parallel_transform_reduce<_RepackedTp,
@@ -174,7 +168,7 @@ auto
174168
__pattern_fill_async(__hetero_tag<_BackendTag> __tag, _ExecutionPolicy&& __exec, _ForwardIterator __first,
175169
_ForwardIterator __last, const _T& __value)
176170
{
177-
return __pattern_walk1_async(
171+
return __pattern_walk1_async<__par_backend_hetero::access_mode::read_write, /*_IsNoInitRequested=*/false>(
178172
__tag, ::std::forward<_ExecutionPolicy>(__exec),
179173
__par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__first),
180174
__par_backend_hetero::make_iter_mode<__par_backend_hetero::access_mode::write>(__last),
@@ -195,9 +189,10 @@ __pattern_transform_scan_base_async(__hetero_tag<_BackendTag>, _ExecutionPolicy&
195189
assert(__first < __last);
196190

197191
auto __n = __last - __first;
198-
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read, _Iterator1>();
192+
auto __keep1 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read>();
199193
auto __buf1 = __keep1(__first, __last);
200-
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write, _Iterator2>();
194+
auto __keep2 = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::write,
195+
/*_IsNoInitRequested=*/true>();
201196
auto __buf2 = __keep2(__result, __result + __n);
202197

203198
auto __res = oneapi::dpl::__par_backend_hetero::__parallel_transform_scan(

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

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -47,7 +47,8 @@ transform_async(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _ForwardIt
4747
const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first, __result);
4848

4949
wait_for_all(::std::forward<_Events>(__dependencies)...);
50-
auto ret_val = oneapi::dpl::__internal::__pattern_walk2_async(
50+
auto ret_val = oneapi::dpl::__internal::__pattern_walk2_async<__par_backend_hetero::access_mode::write,
51+
/*_IsNoInitRequested=*/true>(
5152
__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last, __result,
5253
oneapi::dpl::__internal::__transform_functor<_UnaryOperation>{::std::move(__op)});
5354
return ret_val;
@@ -65,7 +66,8 @@ transform_async(_ExecutionPolicy&& __exec, _ForwardIterator1 __first1, _ForwardI
6566
const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first1, __first2, __result);
6667

6768
wait_for_all(::std::forward<_Events>(__dependencies)...);
68-
auto ret_val = oneapi::dpl::__internal::__pattern_walk3_async(
69+
auto ret_val = oneapi::dpl::__internal::__pattern_walk3_async<__par_backend_hetero::access_mode::write,
70+
/*_IsNoInitRequested=*/true>(
6971
__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first1, __last1, __first2, __result,
7072
oneapi::dpl::__internal::__transform_functor<_BinaryOperation>(::std::move(__op)));
7173
return ret_val;
@@ -97,7 +99,7 @@ sort_async(_ExecutionPolicy&& __exec, _Iterator __first, _Iterator __last, _Comp
9799
wait_for_all(::std::forward<_Events>(__dependencies)...);
98100
assert(__last - __first >= 2);
99101

100-
auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read_write, _Iterator>();
102+
auto __keep = oneapi::dpl::__ranges::__get_sycl_range<__par_backend_hetero::access_mode::read_write>();
101103
auto __buf = __keep(__first, __last);
102104

103105
const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
@@ -128,7 +130,8 @@ for_each_async(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIter
128130
const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
129131

130132
wait_for_all(::std::forward<_Events>(__dependencies)...);
131-
auto ret_val = oneapi::dpl::__internal::__pattern_walk1_async(
133+
auto ret_val = oneapi::dpl::__internal::__pattern_walk1_async<__par_backend_hetero::access_mode::read_write,
134+
/*_IsNoInitRequested=*/false>(
132135
__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first, __last, __f);
133136
return ret_val;
134137
}

‎include/oneapi/dpl/internal/binary_search_impl.h‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -194,13 +194,13 @@ lower_bound_impl(__internal::__hetero_tag<_BackendTag>, Policy&& policy, InputIt
194194

195195
const auto value_size = std::distance(value_start, value_end);
196196

197-
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator1>();
197+
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
198198
auto input_buf = keep_input(start, end);
199199

200-
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator2>();
200+
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
201201
auto value_buf = keep_values(value_start, value_end);
202202

203-
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write, OutputIterator>();
203+
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write>();
204204
auto result_buf = keep_result(result, result + value_size);
205205
auto zip_vw = make_zip_view(input_buf.all_view(), value_buf.all_view(), result_buf.all_view());
206206
const bool use_32bit_indexing = size <= std::numeric_limits<std::uint32_t>::max();
@@ -226,13 +226,13 @@ upper_bound_impl(__internal::__hetero_tag<_BackendTag>, Policy&& policy, InputIt
226226

227227
const auto value_size = std::distance(value_start, value_end);
228228

229-
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator1>();
229+
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
230230
auto input_buf = keep_input(start, end);
231231

232-
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator2>();
232+
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
233233
auto value_buf = keep_values(value_start, value_end);
234234

235-
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write, OutputIterator>();
235+
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write>();
236236
auto result_buf = keep_result(result, result + value_size);
237237
auto zip_vw = make_zip_view(input_buf.all_view(), value_buf.all_view(), result_buf.all_view());
238238
const bool use_32bit_indexing = size <= std::numeric_limits<std::uint32_t>::max();
@@ -258,13 +258,13 @@ binary_search_impl(__internal::__hetero_tag<_BackendTag>, Policy&& policy, Input
258258

259259
const auto value_size = std::distance(value_start, value_end);
260260

261-
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator1>();
261+
auto keep_input = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
262262
auto input_buf = keep_input(start, end);
263263

264-
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read, InputIterator2>();
264+
auto keep_values = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read>();
265265
auto value_buf = keep_values(value_start, value_end);
266266

267-
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write, OutputIterator>();
267+
auto keep_result = oneapi::dpl::__ranges::__get_sycl_range<__bknd::access_mode::read_write>();
268268
auto result_buf = keep_result(result, result + value_size);
269269
auto zip_vw = make_zip_view(input_buf.all_view(), value_buf.all_view(), result_buf.all_view());
270270
const bool use_32bit_indexing = size <= std::numeric_limits<std::uint32_t>::max();

0 commit comments

Comments
 (0)