CUB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups
device_run_length_encode.cuh
Go to the documentation of this file.
1 
2 /******************************************************************************
3  * Copyright (c) 2011, Duane Merrill. All rights reserved.
4  * Copyright (c) 2011-2015, NVIDIA CORPORATION. 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 
35 #pragma once
36 
37 #include <stdio.h>
38 #include <iterator>
39 
40 #include "dispatch/dispatch_rle.cuh"
41 #include "dispatch/dispatch_reduce_by_key.cuh"
42 #include "../util_namespace.cuh"
43 
45 CUB_NS_PREFIX
46 
48 namespace cub {
49 
50 
79 {
80 
143  template <
144  typename InputIteratorT,
145  typename UniqueOutputIteratorT,
146  typename LengthsOutputIteratorT,
147  typename NumRunsOutputIteratorT>
148  CUB_RUNTIME_FUNCTION __forceinline__
149  static cudaError_t Encode(
150  void* d_temp_storage,
151  size_t &temp_storage_bytes,
152  InputIteratorT d_in,
153  UniqueOutputIteratorT d_unique_out,
154  LengthsOutputIteratorT d_counts_out,
155  NumRunsOutputIteratorT d_num_runs_out,
156  int num_items,
157  cudaStream_t stream = 0,
158  bool debug_synchronous = false)
159  {
160  // Data type of value iterator
161  typedef typename std::iterator_traits<LengthsOutputIteratorT>::value_type Value;
162 
163  typedef int OffsetT; // Signed integer type for global offsets
164  typedef NullType* FlagIterator; // FlagT iterator type (not used)
165  typedef NullType SelectOp; // Selection op (not used)
166  typedef Equality EqualityOp; // Default == operator
167  typedef cub::Sum ReductionOp; // Value reduction operator
168 
169  // Generator type for providing 1s values for run-length reduction
170  typedef ConstantInputIterator<Value, OffsetT> LengthsInputIteratorT;
171 
172  Value one_val;
173  one_val = 1;
174 
175  return DispatchReduceByKey<InputIteratorT, UniqueOutputIteratorT, LengthsInputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
176  d_temp_storage,
177  temp_storage_bytes,
178  d_in,
179  d_unique_out,
180  LengthsInputIteratorT(one_val),
181  d_counts_out,
182  d_num_runs_out,
183  EqualityOp(),
184  ReductionOp(),
185  num_items,
186  stream,
187  debug_synchronous);
188  }
189 
190 
241  template <
242  typename InputIteratorT,
243  typename OffsetsOutputIteratorT,
244  typename LengthsOutputIteratorT,
245  typename NumRunsOutputIteratorT>
246  CUB_RUNTIME_FUNCTION __forceinline__
247  static cudaError_t NonTrivialRuns(
248  void* d_temp_storage,
249  size_t &temp_storage_bytes,
250  InputIteratorT d_in,
251  OffsetsOutputIteratorT d_offsets_out,
252  LengthsOutputIteratorT d_lengths_out,
253  NumRunsOutputIteratorT d_num_runs_out,
254  int num_items,
255  cudaStream_t stream = 0,
256  bool debug_synchronous = false)
257  {
258  typedef int OffsetT; // Signed integer type for global offsets
259  typedef Equality EqualityOp; // Default == operator
260 
261  return DeviceRleDispatch<InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, OffsetT>::Dispatch(
262  d_temp_storage,
263  temp_storage_bytes,
264  d_in,
265  d_offsets_out,
266  d_lengths_out,
267  d_num_runs_out,
268  EqualityOp(),
269  num_items,
270  stream,
271  debug_synchronous);
272  }
273 
274 
275 };
276 
277 
278 } // CUB namespace
279 CUB_NS_POSTFIX // Optional outer namespace(s)
280 
281