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 
142  template <
143  typename InputIteratorT,
144  typename UniqueOutputIteratorT,
145  typename LengthsOutputIteratorT,
146  typename NumRunsOutputIteratorT>
147  CUB_RUNTIME_FUNCTION __forceinline__
148  static cudaError_t Encode(
149  void* d_temp_storage,
150  size_t &temp_storage_bytes,
151  InputIteratorT d_in,
152  UniqueOutputIteratorT d_unique_out,
153  LengthsOutputIteratorT d_counts_out,
154  NumRunsOutputIteratorT d_num_runs_out,
155  int num_items,
156  cudaStream_t stream = 0,
157  bool debug_synchronous = false)
158  {
159  // Data type of value iterator
160  typedef typename std::iterator_traits<LengthsOutputIteratorT>::value_type Value;
161 
162  typedef int OffsetT; // Signed integer type for global offsets
163  typedef NullType* FlagIterator; // FlagT iterator type (not used)
164  typedef NullType SelectOp; // Selection op (not used)
165  typedef Equality EqualityOp; // Default == operator
166  typedef cub::Sum ReductionOp; // Value reduction operator
167 
168  // Generator type for providing 1s values for run-length reduction
169  typedef ConstantInputIterator<Value, OffsetT> LengthsInputIteratorT;
170 
171  Value one_val;
172  one_val = 1;
173 
174  return DispatchReduceByKey<InputIteratorT, UniqueOutputIteratorT, LengthsInputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, ReductionOp, OffsetT>::Dispatch(
175  d_temp_storage,
176  temp_storage_bytes,
177  d_in,
178  d_unique_out,
179  LengthsInputIteratorT(one_val),
180  d_counts_out,
181  d_num_runs_out,
182  EqualityOp(),
183  ReductionOp(),
184  num_items,
185  stream,
186  debug_synchronous);
187  }
188 
189 
239  template <
240  typename InputIteratorT,
241  typename OffsetsOutputIteratorT,
242  typename LengthsOutputIteratorT,
243  typename NumRunsOutputIteratorT>
244  CUB_RUNTIME_FUNCTION __forceinline__
245  static cudaError_t NonTrivialRuns(
246  void* d_temp_storage,
247  size_t &temp_storage_bytes,
248  InputIteratorT d_in,
249  OffsetsOutputIteratorT d_offsets_out,
250  LengthsOutputIteratorT d_lengths_out,
251  NumRunsOutputIteratorT d_num_runs_out,
252  int num_items,
253  cudaStream_t stream = 0,
254  bool debug_synchronous = false)
255  {
256  typedef int OffsetT; // Signed integer type for global offsets
257  typedef Equality EqualityOp; // Default == operator
258 
259  return DeviceRleDispatch<InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, NumRunsOutputIteratorT, EqualityOp, OffsetT>::Dispatch(
260  d_temp_storage,
261  temp_storage_bytes,
262  d_in,
263  d_offsets_out,
264  d_lengths_out,
265  d_num_runs_out,
266  EqualityOp(),
267  num_items,
268  stream,
269  debug_synchronous);
270  }
271 
272 
273 };
274 
275 
276 } // CUB namespace
277 CUB_NS_POSTFIX // Optional outer namespace(s)
278 
279