/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.0/hipcub/include/hipcub/backend/rocprim/device/device_scan.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.0/hipcub/include/hipcub/backend/rocprim/device/device_scan.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.0/hipcub/include/hipcub/backend/rocprim/device/device_scan.hpp Source File
device_scan.hpp
1 /******************************************************************************
2  * Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
4  * Modifications Copyright (c) 2017-2020, Advanced Micro Devices, Inc. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  *
28  ******************************************************************************/
29 
30 #ifndef HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_
31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_
32 
33 #include <iostream>
34 #include "../../../config.hpp"
35 
36 #include "../thread/thread_operators.hpp"
37 
38 #include <rocprim/device/device_scan.hpp>
39 #include <rocprim/device/device_scan_by_key.hpp>
40 
41 BEGIN_HIPCUB_NAMESPACE
42 
43 class DeviceScan
44 {
45 public:
46  template <
47  typename InputIteratorT,
48  typename OutputIteratorT
49  >
50  HIPCUB_RUNTIME_FUNCTION static
51  hipError_t InclusiveSum(void *d_temp_storage,
52  size_t &temp_storage_bytes,
53  InputIteratorT d_in,
54  OutputIteratorT d_out,
55  size_t num_items,
56  hipStream_t stream = 0,
57  bool debug_synchronous = false)
58  {
59  return InclusiveScan(
60  d_temp_storage, temp_storage_bytes,
61  d_in, d_out, ::hipcub::Sum(), num_items,
62  stream, debug_synchronous
63  );
64  }
65 
66  template <
67  typename InputIteratorT,
68  typename OutputIteratorT,
69  typename ScanOpT
70  >
71  HIPCUB_RUNTIME_FUNCTION static
72  hipError_t InclusiveScan(void *d_temp_storage,
73  size_t &temp_storage_bytes,
74  InputIteratorT d_in,
75  OutputIteratorT d_out,
76  ScanOpT scan_op,
77  size_t num_items,
78  hipStream_t stream = 0,
79  bool debug_synchronous = false)
80  {
81  return ::rocprim::inclusive_scan(
82  d_temp_storage, temp_storage_bytes,
83  d_in, d_out, num_items,
84  scan_op,
85  stream, debug_synchronous
86  );
87  }
88 
89  template <
90  typename InputIteratorT,
91  typename OutputIteratorT
92  >
93  HIPCUB_RUNTIME_FUNCTION static
94  hipError_t ExclusiveSum(void *d_temp_storage,
95  size_t &temp_storage_bytes,
96  InputIteratorT d_in,
97  OutputIteratorT d_out,
98  size_t num_items,
99  hipStream_t stream = 0,
100  bool debug_synchronous = false)
101  {
102  using T = typename std::iterator_traits<InputIteratorT>::value_type;
103  return ExclusiveScan(
104  d_temp_storage, temp_storage_bytes,
105  d_in, d_out, ::hipcub::Sum(), T(0), num_items,
106  stream, debug_synchronous
107  );
108  }
109 
110  template <
111  typename InputIteratorT,
112  typename OutputIteratorT,
113  typename ScanOpT,
114  typename InitValueT
115  >
116  HIPCUB_RUNTIME_FUNCTION static
117  hipError_t ExclusiveScan(void *d_temp_storage,
118  size_t &temp_storage_bytes,
119  InputIteratorT d_in,
120  OutputIteratorT d_out,
121  ScanOpT scan_op,
122  InitValueT init_value,
123  size_t num_items,
124  hipStream_t stream = 0,
125  bool debug_synchronous = false)
126  {
127  return ::rocprim::exclusive_scan(
128  d_temp_storage, temp_storage_bytes,
129  d_in, d_out, init_value, num_items,
130  scan_op,
131  stream, debug_synchronous
132  );
133  }
134 
135  template <
136  typename InputIteratorT,
137  typename OutputIteratorT,
138  typename ScanOpT,
139  typename InitValueT,
140  typename InitValueIterT = InitValueT*
141  >
142  HIPCUB_RUNTIME_FUNCTION static
143  hipError_t ExclusiveScan(void *d_temp_storage,
144  size_t &temp_storage_bytes,
145  InputIteratorT d_in,
146  OutputIteratorT d_out,
147  ScanOpT scan_op,
148  FutureValue<InitValueT, InitValueIterT> init_value,
149  int num_items,
150  hipStream_t stream = 0,
151  bool debug_synchronous = false)
152  {
153  return ::rocprim::exclusive_scan(
154  d_temp_storage, temp_storage_bytes,
155  d_in, d_out, init_value, num_items,
156  scan_op,
157  stream, debug_synchronous
158  );
159  }
160 
161  template <
162  typename KeysInputIteratorT,
163  typename ValuesInputIteratorT,
164  typename ValuesOutputIteratorT,
165  typename EqualityOpT = ::hipcub::Equality
166  >
167  HIPCUB_RUNTIME_FUNCTION static
168  hipError_t ExclusiveSumByKey(void *d_temp_storage,
169  size_t &temp_storage_bytes,
170  KeysInputIteratorT d_keys_in,
171  ValuesInputIteratorT d_values_in,
172  ValuesOutputIteratorT d_values_out,
173  int num_items,
174  EqualityOpT equality_op = EqualityOpT(),
175  hipStream_t stream = 0,
176  bool debug_synchronous = false)
177  {
178  using in_value_type = typename std::iterator_traits<ValuesInputIteratorT>::value_type;
179 
180  return ::rocprim::exclusive_scan_by_key(
181  d_temp_storage, temp_storage_bytes,
182  d_keys_in, d_values_in, d_values_out,
183  static_cast<in_value_type>(0), static_cast<size_t>(num_items),
184  ::hipcub::Sum(), equality_op, stream, debug_synchronous
185  );
186  }
187 
188  template <
189  typename KeysInputIteratorT,
190  typename ValuesInputIteratorT,
191  typename ValuesOutputIteratorT,
192  typename ScanOpT,
193  typename InitValueT,
194  typename EqualityOpT = ::hipcub::Equality
195  >
196  HIPCUB_RUNTIME_FUNCTION static
197  hipError_t ExclusiveScanByKey(void *d_temp_storage,
198  size_t &temp_storage_bytes,
199  KeysInputIteratorT d_keys_in,
200  ValuesInputIteratorT d_values_in,
201  ValuesOutputIteratorT d_values_out,
202  ScanOpT scan_op,
203  InitValueT init_value,
204  int num_items,
205  EqualityOpT equality_op = EqualityOpT(),
206  hipStream_t stream = 0,
207  bool debug_synchronous = false)
208  {
209  return ::rocprim::exclusive_scan_by_key(
210  d_temp_storage, temp_storage_bytes,
211  d_keys_in, d_values_in, d_values_out,
212  init_value, static_cast<size_t>(num_items),
213  scan_op, equality_op, stream, debug_synchronous
214  );
215  }
216 
217  template <
218  typename KeysInputIteratorT,
219  typename ValuesInputIteratorT,
220  typename ValuesOutputIteratorT,
221  typename EqualityOpT = ::hipcub::Equality
222  >
223  HIPCUB_RUNTIME_FUNCTION static
224  hipError_t InclusiveSumByKey(void *d_temp_storage,
225  size_t &temp_storage_bytes,
226  KeysInputIteratorT d_keys_in,
227  ValuesInputIteratorT d_values_in,
228  ValuesOutputIteratorT d_values_out,
229  int num_items,
230  EqualityOpT equality_op = EqualityOpT(),
231  hipStream_t stream = 0,
232  bool debug_synchronous = false)
233  {
234  return ::rocprim::inclusive_scan_by_key(
235  d_temp_storage, temp_storage_bytes,
236  d_keys_in, d_values_in, d_values_out,
237  static_cast<size_t>(num_items), ::hipcub::Sum(),
238  equality_op, stream, debug_synchronous
239  );
240  }
241 
242  template <
243  typename KeysInputIteratorT,
244  typename ValuesInputIteratorT,
245  typename ValuesOutputIteratorT,
246  typename ScanOpT,
247  typename EqualityOpT = ::hipcub::Equality
248  >
249  HIPCUB_RUNTIME_FUNCTION static
250  hipError_t InclusiveScanByKey(void *d_temp_storage,
251  size_t &temp_storage_bytes,
252  KeysInputIteratorT d_keys_in,
253  ValuesInputIteratorT d_values_in,
254  ValuesOutputIteratorT d_values_out,
255  ScanOpT scan_op,
256  int num_items,
257  EqualityOpT equality_op = EqualityOpT(),
258  hipStream_t stream = 0,
259  bool debug_synchronous = false)
260  {
261  return ::rocprim::inclusive_scan_by_key(
262  d_temp_storage, temp_storage_bytes,
263  d_keys_in, d_values_in, d_values_out,
264  static_cast<size_t>(num_items), scan_op,
265  equality_op, stream, debug_synchronous
266  );
267  }
268 };
269 
270 END_HIPCUB_NAMESPACE
271 
272 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_
Definition: thread_operators.hpp:40
Definition: thread_operators.hpp:76