RAJA
RAJA provides a collection of platform portability abstractions for C++ HPC applications.
sort.hpp
Go to the documentation of this file.
1 
11 //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
12 // Copyright (c) Lawrence Livermore National Security, LLC and other
13 // RAJA Project Developers. See top-level LICENSE and COPYRIGHT
14 // files for dates and other details. No copyright assignment is required
15 // to contribute to RAJA.
16 //
17 // SPDX-License-Identifier: (BSD-3-Clause)
18 //~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~//
19 
20 #ifndef RAJA_sort_hip_HPP
21 #define RAJA_sort_hip_HPP
22 
23 #include "RAJA/config.hpp"
24 
25 #if defined(RAJA_ENABLE_HIP)
26 
27 #include <climits>
28 #include <iterator>
29 #include <type_traits>
30 
31 #if defined(__HIPCC__)
32 // Tell rocprim to provide its HIP API
33 #define ROCPRIM_HIP_API 1
34 #include "rocprim/device/device_transform.hpp"
35 #include "rocprim/device/device_radix_sort.hpp"
36 #elif defined(__CUDACC__)
37 #include "cub/device/device_radix_sort.cuh"
38 #endif
39 
40 #include "RAJA/util/concepts.hpp"
41 #include "RAJA/util/Operators.hpp"
42 #include "RAJA/pattern/detail/algorithm.hpp"
43 #include "RAJA/policy/hip/MemUtils_HIP.hpp"
44 #include "RAJA/policy/hip/policy.hpp"
45 
46 namespace RAJA
47 {
48 namespace impl
49 {
50 namespace sort
51 {
52 
53 namespace detail
54 {
55 
56 #if defined(__HIPCC__)
57 template<typename R>
58 using double_buffer = ::rocprim::double_buffer<R>;
59 #elif defined(__CUDACC__)
60 template<typename R>
61 using double_buffer = ::cub::DoubleBuffer<R>;
62 #endif
63 
64 template<typename R>
65 R* get_current(double_buffer<R>& d_bufs)
66 {
67 #if defined(__HIPCC__)
68  return d_bufs.current();
69 #elif defined(__CUDACC__)
70  return d_bufs.Current();
71 #endif
72 }
73 
74 } // namespace detail
75 
79 template<typename IterationMapping,
80  typename IterationGetter,
81  typename Concretizer,
82  bool Async,
83  typename Iter,
84  typename Compare>
85 concepts::enable_if_t<
86  resources::EventProxy<resources::Hip>,
87  concepts::negate<concepts::all_of<
88  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
89  std::is_pointer<Iter>,
90  concepts::any_of<
91  camp::is_same<Compare,
92  operators::less<RAJA::detail::IterVal<Iter>>>,
93  camp::is_same<Compare,
94  operators::greater<RAJA::detail::IterVal<Iter>>>>>>>
95 stable(resources::Hip hip_res,
96  ::RAJA::policy::hip::
97  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
98  Iter,
99  Iter,
100  Compare)
101 {
102  static_assert(
103  concepts::all_of<
104  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
105  std::is_pointer<Iter>,
106  concepts::any_of<
107  camp::is_same<Compare,
108  operators::less<RAJA::detail::IterVal<Iter>>>,
109  camp::is_same<Compare, operators::greater<
110  RAJA::detail::IterVal<Iter>>>>>::value,
111  "RAJA stable_sort<hip_exec> is only implemented for pointers to "
112  "arithmetic types and RAJA::operators::less and "
113  "RAJA::operators::greater.");
114 
115  return resources::EventProxy<resources::Hip>(hip_res);
116 }
117 
121 template<typename IterationMapping,
122  typename IterationGetter,
123  typename Concretizer,
124  bool Async,
125  typename Iter>
126 concepts::enable_if_t<resources::EventProxy<resources::Hip>,
127  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
128  std::is_pointer<Iter>>
129 stable(resources::Hip hip_res,
130  ::RAJA::policy::hip::
131  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
132  Iter begin,
133  Iter end,
134  operators::less<RAJA::detail::IterVal<Iter>>)
135 {
136  hipStream_t stream = hip_res.get_stream();
137 
138  using R = RAJA::detail::IterVal<Iter>;
139 
140  int len = std::distance(begin, end);
141  int begin_bit = 0;
142  int end_bit = sizeof(R) * CHAR_BIT;
143 
144  // Allocate temporary storage for the output array
145  R* d_out = hip::device_mempool_type::getInstance().malloc<R>(len);
146 
147  // use cub double buffer to reduce temporary memory requirements
148  // by allowing cub to write to the begin buffer
149  detail::double_buffer<R> d_keys(begin, d_out);
150 
151  // Determine temporary device storage requirements
152  void* d_temp_storage = nullptr;
153  size_t temp_storage_bytes = 0;
154 #if defined(__HIPCC__)
155  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_keys, d_temp_storage,
156  temp_storage_bytes, d_keys, len, begin_bit,
157  end_bit, stream);
158 #elif defined(__CUDACC__)
159  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortKeys,
160  d_temp_storage, temp_storage_bytes, d_keys,
161  len, begin_bit, end_bit, stream);
162 #endif
163  // Allocate temporary storage
164  d_temp_storage =
165  hip::device_mempool_type::getInstance().malloc<unsigned char>(
166  temp_storage_bytes);
167 
168  // Run
169 #if defined(__HIPCC__)
170  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_keys, d_temp_storage,
171  temp_storage_bytes, d_keys, len, begin_bit,
172  end_bit, stream);
173 #elif defined(__CUDACC__)
174  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortKeys,
175  d_temp_storage, temp_storage_bytes, d_keys,
176  len, begin_bit, end_bit, stream);
177 #endif
178  // Free temporary storage
179  hip::device_mempool_type::getInstance().free(d_temp_storage);
180 
181  if (detail::get_current(d_keys) == d_out)
182  {
183 
184  // copy
185  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, begin, d_out, len * sizeof(R),
186  hipMemcpyDefault, stream);
187  }
188 
189  hip::device_mempool_type::getInstance().free(d_out);
190 
191  hip::launch(hip_res, Async);
192 
193  return resources::EventProxy<resources::Hip>(hip_res);
194 }
195 
199 template<typename IterationMapping,
200  typename IterationGetter,
201  typename Concretizer,
202  bool Async,
203  typename Iter>
204 concepts::enable_if_t<resources::EventProxy<resources::Hip>,
205  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
206  std::is_pointer<Iter>>
207 stable(resources::Hip hip_res,
208  ::RAJA::policy::hip::
209  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
210  Iter begin,
211  Iter end,
212  operators::greater<RAJA::detail::IterVal<Iter>>)
213 {
214  hipStream_t stream = hip_res.get_stream();
215 
216  using R = RAJA::detail::IterVal<Iter>;
217 
218  int len = std::distance(begin, end);
219  int begin_bit = 0;
220  int end_bit = sizeof(R) * CHAR_BIT;
221 
222  // Allocate temporary storage for the output array
223  R* d_out = hip::device_mempool_type::getInstance().malloc<R>(len);
224 
225  // use cub double buffer to reduce temporary memory requirements
226  // by allowing cub to write to the begin buffer
227  detail::double_buffer<R> d_keys(begin, d_out);
228 
229  // Determine temporary device storage requirements
230  void* d_temp_storage = nullptr;
231  size_t temp_storage_bytes = 0;
232 #if defined(__HIPCC__)
233  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_keys_desc, d_temp_storage,
234  temp_storage_bytes, d_keys, len, begin_bit,
235  end_bit, stream);
236 #elif defined(__CUDACC__)
237  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortKeysDescending,
238  d_temp_storage, temp_storage_bytes, d_keys,
239  len, begin_bit, end_bit, stream);
240 #endif
241  // Allocate temporary storage
242  d_temp_storage =
243  hip::device_mempool_type::getInstance().malloc<unsigned char>(
244  temp_storage_bytes);
245 
246  // Run
247 #if defined(__HIPCC__)
248  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_keys_desc, d_temp_storage,
249  temp_storage_bytes, d_keys, len, begin_bit,
250  end_bit, stream);
251 #elif defined(__CUDACC__)
252  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortKeysDescending,
253  d_temp_storage, temp_storage_bytes, d_keys,
254  len, begin_bit, end_bit, stream);
255 #endif
256  // Free temporary storage
257  hip::device_mempool_type::getInstance().free(d_temp_storage);
258 
259  if (detail::get_current(d_keys) == d_out)
260  {
261 
262  // copy
263  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, begin, d_out, len * sizeof(R),
264  hipMemcpyDefault, stream);
265  }
266 
267  hip::device_mempool_type::getInstance().free(d_out);
268 
269  hip::launch(hip_res, Async);
270 
271  return resources::EventProxy<resources::Hip>(hip_res);
272 }
273 
277 template<typename IterationMapping,
278  typename IterationGetter,
279  typename Concretizer,
280  bool Async,
281  typename Iter,
282  typename Compare>
283 concepts::enable_if_t<
284  resources::EventProxy<resources::Hip>,
285  concepts::negate<concepts::all_of<
286  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
287  std::is_pointer<Iter>,
288  concepts::any_of<
289  camp::is_same<Compare,
290  operators::less<RAJA::detail::IterVal<Iter>>>,
291  camp::is_same<Compare,
292  operators::greater<RAJA::detail::IterVal<Iter>>>>>>>
293 unstable(resources::Hip hip_res,
294  ::RAJA::policy::hip::
295  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
296  Iter,
297  Iter,
298  Compare)
299 {
300  static_assert(
301  concepts::all_of<
302  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
303  std::is_pointer<Iter>,
304  concepts::any_of<
305  camp::is_same<Compare,
306  operators::less<RAJA::detail::IterVal<Iter>>>,
307  camp::is_same<Compare, operators::greater<
308  RAJA::detail::IterVal<Iter>>>>>::value,
309  "RAJA sort<hip_exec> is only implemented for pointers to arithmetic "
310  "types and RAJA::operators::less and RAJA::operators::greater.");
311 
312  return resources::EventProxy<resources::Hip>(hip_res);
313 }
314 
318 template<typename IterationMapping,
319  typename IterationGetter,
320  typename Concretizer,
321  bool Async,
322  typename Iter>
323 concepts::enable_if_t<resources::EventProxy<resources::Hip>,
324  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
325  std::is_pointer<Iter>>
326 unstable(resources::Hip hip_res,
327  ::RAJA::policy::hip::
328  hip_exec<IterationMapping, IterationGetter, Concretizer, Async> p,
329  Iter begin,
330  Iter end,
331  operators::less<RAJA::detail::IterVal<Iter>> comp)
332 {
333  return stable(hip_res, p, begin, end, comp);
334 }
335 
339 template<typename IterationMapping,
340  typename IterationGetter,
341  typename Concretizer,
342  bool Async,
343  typename Iter>
344 concepts::enable_if_t<resources::EventProxy<resources::Hip>,
345  type_traits::is_arithmetic<RAJA::detail::IterVal<Iter>>,
346  std::is_pointer<Iter>>
347 unstable(resources::Hip hip_res,
348  ::RAJA::policy::hip::
349  hip_exec<IterationMapping, IterationGetter, Concretizer, Async> p,
350  Iter begin,
351  Iter end,
352  operators::greater<RAJA::detail::IterVal<Iter>> comp)
353 {
354  return stable(hip_res, p, begin, end, comp);
355 }
356 
360 template<typename IterationMapping,
361  typename IterationGetter,
362  typename Concretizer,
363  bool Async,
364  typename KeyIter,
365  typename ValIter,
366  typename Compare>
367 concepts::enable_if_t<
368  resources::EventProxy<resources::Hip>,
369  concepts::negate<concepts::all_of<
370  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
371  std::is_pointer<KeyIter>,
372  std::is_pointer<ValIter>,
373  concepts::any_of<
374  camp::is_same<Compare,
375  operators::less<RAJA::detail::IterVal<KeyIter>>>,
376  camp::is_same<
377  Compare,
378  operators::greater<RAJA::detail::IterVal<KeyIter>>>>>>>
379 stable_pairs(
380  resources::Hip hip_res,
381  ::RAJA::policy::hip::
382  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
383  KeyIter,
384  KeyIter,
385  ValIter,
386  Compare)
387 {
388  static_assert(std::is_pointer<KeyIter>::value,
389  "stable_sort_pairs<hip_exec> is only implemented for pointers");
390  static_assert(std::is_pointer<ValIter>::value,
391  "stable_sort_pairs<hip_exec> is only implemented for pointers");
392  using K = RAJA::detail::IterVal<KeyIter>;
393  static_assert(
394  type_traits::is_arithmetic<K>::value,
395  "stable_sort_pairs<hip_exec> is only implemented for arithmetic types");
396  static_assert(
397  concepts::any_of<camp::is_same<Compare, operators::less<K>>,
398  camp::is_same<Compare, operators::greater<K>>>::value,
399  "stable_sort_pairs<hip_exec> is only implemented for "
400  "RAJA::operators::less or RAJA::operators::greater");
401 
402  return resources::EventProxy<resources::Hip>(hip_res);
403 }
404 
408 template<typename IterationMapping,
409  typename IterationGetter,
410  typename Concretizer,
411  bool Async,
412  typename KeyIter,
413  typename ValIter>
414 concepts::enable_if_t<
415  resources::EventProxy<resources::Hip>,
416  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
417  std::is_pointer<KeyIter>,
418  std::is_pointer<ValIter>>
419 stable_pairs(
420  resources::Hip hip_res,
421  ::RAJA::policy::hip::
422  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
423  KeyIter keys_begin,
424  KeyIter keys_end,
425  ValIter vals_begin,
426  operators::less<RAJA::detail::IterVal<KeyIter>>)
427 {
428  hipStream_t stream = hip_res.get_stream();
429 
430  using K = RAJA::detail::IterVal<KeyIter>;
431  using V = RAJA::detail::IterVal<ValIter>;
432 
433  int len = std::distance(keys_begin, keys_end);
434  int begin_bit = 0;
435  int end_bit = sizeof(K) * CHAR_BIT;
436 
437  // Allocate temporary storage for the output arrays
438  K* d_keys_out = hip::device_mempool_type::getInstance().malloc<K>(len);
439  V* d_vals_out = hip::device_mempool_type::getInstance().malloc<V>(len);
440 
441  // use cub double buffer to reduce temporary memory requirements
442  // by allowing cub to write to the keys_begin and vals_begin buffers
443  detail::double_buffer<K> d_keys(keys_begin, d_keys_out);
444  detail::double_buffer<V> d_vals(vals_begin, d_vals_out);
445 
446  // Determine temporary device storage requirements
447  void* d_temp_storage = nullptr;
448  size_t temp_storage_bytes = 0;
449 #if defined(__HIPCC__)
450  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_pairs, d_temp_storage,
451  temp_storage_bytes, d_keys, d_vals, len,
452  begin_bit, end_bit, stream);
453 #elif defined(__CUDACC__)
454  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortPairs,
455  d_temp_storage, temp_storage_bytes, d_keys,
456  d_vals, len, begin_bit, end_bit, stream);
457 #endif
458  // Allocate temporary storage
459  d_temp_storage =
460  hip::device_mempool_type::getInstance().malloc<unsigned char>(
461  temp_storage_bytes);
462 
463  // Run
464 #if defined(__HIPCC__)
465  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_pairs, d_temp_storage,
466  temp_storage_bytes, d_keys, d_vals, len,
467  begin_bit, end_bit, stream);
468 #elif defined(__CUDACC__)
469  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortPairs,
470  d_temp_storage, temp_storage_bytes, d_keys,
471  d_vals, len, begin_bit, end_bit, stream);
472 #endif
473  // Free temporary storage
474  hip::device_mempool_type::getInstance().free(d_temp_storage);
475 
476  if (detail::get_current(d_keys) == d_keys_out)
477  {
478 
479  // copy keys
480  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, keys_begin, d_keys_out,
481  len * sizeof(K), hipMemcpyDefault, stream);
482  }
483  if (detail::get_current(d_vals) == d_vals_out)
484  {
485 
486  // copy vals
487  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, vals_begin, d_vals_out,
488  len * sizeof(V), hipMemcpyDefault, stream);
489  }
490 
491  hip::device_mempool_type::getInstance().free(d_keys_out);
492  hip::device_mempool_type::getInstance().free(d_vals_out);
493 
494  hip::launch(hip_res, Async);
495 
496  return resources::EventProxy<resources::Hip>(hip_res);
497 }
498 
502 template<typename IterationMapping,
503  typename IterationGetter,
504  typename Concretizer,
505  bool Async,
506  typename KeyIter,
507  typename ValIter>
508 concepts::enable_if_t<
509  resources::EventProxy<resources::Hip>,
510  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
511  std::is_pointer<KeyIter>,
512  std::is_pointer<ValIter>>
513 stable_pairs(
514  resources::Hip hip_res,
515  ::RAJA::policy::hip::
516  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
517  KeyIter keys_begin,
518  KeyIter keys_end,
519  ValIter vals_begin,
520  operators::greater<RAJA::detail::IterVal<KeyIter>>)
521 {
522  hipStream_t stream = hip_res.get_stream();
523 
524  using K = RAJA::detail::IterVal<KeyIter>;
525  using V = RAJA::detail::IterVal<ValIter>;
526 
527  int len = std::distance(keys_begin, keys_end);
528  int begin_bit = 0;
529  int end_bit = sizeof(K) * CHAR_BIT;
530 
531  // Allocate temporary storage for the output arrays
532  K* d_keys_out = hip::device_mempool_type::getInstance().malloc<K>(len);
533  V* d_vals_out = hip::device_mempool_type::getInstance().malloc<V>(len);
534 
535  // use cub double buffer to reduce temporary memory requirements
536  // by allowing cub to write to the keys_begin and vals_begin buffers
537  detail::double_buffer<K> d_keys(keys_begin, d_keys_out);
538  detail::double_buffer<V> d_vals(vals_begin, d_vals_out);
539 
540  // Determine temporary device storage requirements
541  void* d_temp_storage = nullptr;
542  size_t temp_storage_bytes = 0;
543 #if defined(__HIPCC__)
544  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_pairs_desc,
545  d_temp_storage, temp_storage_bytes, d_keys,
546  d_vals, len, begin_bit, end_bit, stream);
547 #elif defined(__CUDACC__)
548  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortPairsDescending,
549  d_temp_storage, temp_storage_bytes, d_keys,
550  d_vals, len, begin_bit, end_bit, stream);
551 #endif
552  // Allocate temporary storage
553  d_temp_storage =
554  hip::device_mempool_type::getInstance().malloc<unsigned char>(
555  temp_storage_bytes);
556 
557  // Run
558 #if defined(__HIPCC__)
559  CAMP_HIP_API_INVOKE_AND_CHECK(::rocprim::radix_sort_pairs_desc,
560  d_temp_storage, temp_storage_bytes, d_keys,
561  d_vals, len, begin_bit, end_bit, stream);
562 #elif defined(__CUDACC__)
563  CAMP_CUDA_API_INVOKE_AND_CHECK(::cub::DeviceRadixSort::SortPairsDescending,
564  d_temp_storage, temp_storage_bytes, d_keys,
565  d_vals, len, begin_bit, end_bit, stream);
566 #endif
567  // Free temporary storage
568  hip::device_mempool_type::getInstance().free(d_temp_storage);
569 
570  if (detail::get_current(d_keys) == d_keys_out)
571  {
572 
573  // copy keys
574  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, keys_begin, d_keys_out,
575  len * sizeof(K), hipMemcpyDefault, stream);
576  }
577  if (detail::get_current(d_vals) == d_vals_out)
578  {
579 
580  // copy vals
581  CAMP_HIP_API_INVOKE_AND_CHECK(hipMemcpyAsync, vals_begin, d_vals_out,
582  len * sizeof(V), hipMemcpyDefault, stream);
583  }
584 
585  hip::device_mempool_type::getInstance().free(d_keys_out);
586  hip::device_mempool_type::getInstance().free(d_vals_out);
587 
588  hip::launch(hip_res, Async);
589 
590  return resources::EventProxy<resources::Hip>(hip_res);
591 }
592 
596 template<typename IterationMapping,
597  typename IterationGetter,
598  typename Concretizer,
599  bool Async,
600  typename KeyIter,
601  typename ValIter,
602  typename Compare>
603 concepts::enable_if_t<
604  resources::EventProxy<resources::Hip>,
605  concepts::negate<concepts::all_of<
606  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
607  std::is_pointer<KeyIter>,
608  std::is_pointer<ValIter>,
609  concepts::any_of<
610  camp::is_same<Compare,
611  operators::less<RAJA::detail::IterVal<KeyIter>>>,
612  camp::is_same<
613  Compare,
614  operators::greater<RAJA::detail::IterVal<KeyIter>>>>>>>
615 unstable_pairs(
616  resources::Hip hip_res,
617  ::RAJA::policy::hip::
618  hip_exec<IterationMapping, IterationGetter, Concretizer, Async>,
619  KeyIter,
620  KeyIter,
621  ValIter,
622  Compare)
623 {
624  static_assert(std::is_pointer<KeyIter>::value,
625  "sort_pairs<hip_exec> is only implemented for pointers");
626  static_assert(std::is_pointer<ValIter>::value,
627  "sort_pairs<hip_exec> is only implemented for pointers");
628  using K = RAJA::detail::IterVal<KeyIter>;
629  static_assert(
630  type_traits::is_arithmetic<K>::value,
631  "sort_pairs<hip_exec> is only implemented for arithmetic types");
632  static_assert(
633  concepts::any_of<camp::is_same<Compare, operators::less<K>>,
634  camp::is_same<Compare, operators::greater<K>>>::value,
635  "sort_pairs<hip_exec> is only implemented for RAJA::operators::less or "
636  "RAJA::operators::greater");
637 
638  return resources::EventProxy<resources::Hip>(hip_res);
639 }
640 
644 template<typename IterationMapping,
645  typename IterationGetter,
646  typename Concretizer,
647  bool Async,
648  typename KeyIter,
649  typename ValIter>
650 concepts::enable_if_t<
651  resources::EventProxy<resources::Hip>,
652  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
653  std::is_pointer<KeyIter>,
654  std::is_pointer<ValIter>>
655 unstable_pairs(
656  resources::Hip hip_res,
657  ::RAJA::policy::hip::
658  hip_exec<IterationMapping, IterationGetter, Concretizer, Async> p,
659  KeyIter keys_begin,
660  KeyIter keys_end,
661  ValIter vals_begin,
662  operators::less<RAJA::detail::IterVal<KeyIter>> comp)
663 {
664  return stable_pairs(hip_res, p, keys_begin, keys_end, vals_begin, comp);
665 }
666 
670 template<typename IterationMapping,
671  typename IterationGetter,
672  typename Concretizer,
673  bool Async,
674  typename KeyIter,
675  typename ValIter>
676 concepts::enable_if_t<
677  resources::EventProxy<resources::Hip>,
678  type_traits::is_arithmetic<RAJA::detail::IterVal<KeyIter>>,
679  std::is_pointer<KeyIter>,
680  std::is_pointer<ValIter>>
681 unstable_pairs(
682  resources::Hip hip_res,
683  ::RAJA::policy::hip::
684  hip_exec<IterationMapping, IterationGetter, Concretizer, Async> p,
685  KeyIter keys_begin,
686  KeyIter keys_end,
687  ValIter vals_begin,
688  operators::greater<RAJA::detail::IterVal<KeyIter>> comp)
689 {
690  return stable_pairs(hip_res, p, keys_begin, keys_end, vals_begin, comp);
691 }
692 
693 } // namespace sort
694 
695 } // namespace impl
696 
697 } // namespace RAJA
698 
699 #endif // closing endif for RAJA_ENABLE_HIP guard
700 
701 #endif // closing endif for header file include guard
MemUtils_HIP.hpp
Header file defining prototypes for routines used to manage memory for HIP reductions and other opera...
Operators.hpp
Header file for RAJA operator definitions.
algorithm.hpp
Header file for RAJA algorithm definitions.
concepts.hpp
Header file for RAJA concept definitions.
policy.hpp
Header file containing RAJA HIP policy definitions.
RAJA::detail::IterVal
typename ::std::iterator_traits< Iter >::value_type IterVal
Definition: algorithm.hpp:38
RAJA::impl::sort::stable_pairs
concepts::enable_if_t< resources::EventProxy< resources::Host >, type_traits::is_openmp_policy< ExecPolicy > > stable_pairs(resources::Host host_res, const ExecPolicy &, KeyIter keys_begin, KeyIter keys_end, ValIter vals_begin, Compare comp)
Definition: sort.hpp:276
RAJA::impl::sort::stable
concepts::enable_if_t< resources::EventProxy< resources::Host >, type_traits::is_openmp_policy< ExecPolicy > > stable(resources::Host host_res, const ExecPolicy &, Iter begin, Iter end, Compare comp)
stable sort given range using comparison function
Definition: sort.hpp:230
RAJA::impl::sort::unstable
concepts::enable_if_t< resources::EventProxy< resources::Host >, type_traits::is_openmp_policy< ExecPolicy > > unstable(resources::Host host_res, const ExecPolicy &, Iter begin, Iter end, Compare comp)
sort given range using comparison function
Definition: sort.hpp:213
RAJA::impl::sort::unstable_pairs
concepts::enable_if_t< resources::EventProxy< resources::Host >, type_traits::is_openmp_policy< ExecPolicy > > unstable_pairs(resources::Host host_res, const ExecPolicy &, KeyIter keys_begin, KeyIter keys_end, ValIter vals_begin, Compare comp)
sort given range of pairs using comparison function on keys
Definition: sort.hpp:250
RAJA::policy_by_value_interface::sort
concepts::enable_if_t< resources::EventProxy< Res >, type_traits::is_execution_policy< ExecPolicy >, type_traits::is_resource< Res >, std::is_constructible< camp::resources::Resource, Res >, type_traits::is_range< Container > > sort(ExecPolicy &&p, Res r, Container &&c, Compare comp=Compare {})
sort execution pattern
Definition: sort.hpp:61
RAJA
Definition: AlignedRangeIndexSetBuilders.cpp:35
RAJA::launch
void launch(LaunchParams const &launch_params, ReduceParams &&... rest_of_launch_args)
Definition: launch_core.hpp:268
Generated by doxygen 1.9.1