← Back to branch summary

~ubuntu-branches/debian/sid/lammps/sid

« back to all changes in this revision

Viewing changes to lib/kokkos/TPL/cub/warp/warp_reduce.cuh

  • Committer: Package Import Robot
  • Author(s): Anton Gladky
  • Date: 2015-04-29 23:44:49 UTC
  • mfrom: (5.1.3 experimental)
  • Revision ID: package-import@ubuntu.com-20150429234449-mbhy9utku6hp6oq8
Tags: 0~20150313.gitfa668e1-1
Upload into unstable.

Show diffs side-by-side

added added

removed removed

Lines of Context:
lib/kokkos/TPL/cub/warp/warp_reduce.cuh
 
1
/******************************************************************************
 
2
 * Copyright (c) 2011, Duane Merrill.  All rights reserved.
 
3
 * Copyright (c) 2011-2013, NVIDIA CORPORATION.  All rights reserved.
 
4
 * 
 
5
 * Redistribution and use in source and binary forms, with or without
 
6
 * modification, are permitted provided that the following conditions are met:
 
7
 *     * Redistributions of source code must retain the above copyright
 
8
 *       notice, this list of conditions and the following disclaimer.
 
9
 *     * Redistributions in binary form must reproduce the above copyright
 
10
 *       notice, this list of conditions and the following disclaimer in the
 
11
 *       documentation and/or other materials provided with the distribution.
 
12
 *     * Neither the name of the NVIDIA CORPORATION nor the
 
13
 *       names of its contributors may be used to endorse or promote products
 
14
 *       derived from this software without specific prior written permission.
 
15
 * 
 
16
 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
 
17
 * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 
18
 * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
 
19
 * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
 
20
 * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
 
21
 * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
 
22
 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
 
23
 * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 
24
 * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
 
25
 * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
26
 *
 
27
 ******************************************************************************/
 
28
 
 
29
/**
 
30
 * \file
 
31
 * The cub::WarpReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across CUDA warp threads.
 
32
 */
 
33
 
 
34
#pragma once
 
35
 
 
36
#include "specializations/warp_reduce_shfl.cuh"
 
37
#include "specializations/warp_reduce_smem.cuh"
 
38
#include "../thread/thread_operators.cuh"
 
39
#include "../util_arch.cuh"
 
40
#include "../util_type.cuh"
 
41
#include "../util_namespace.cuh"
 
42
 
 
43
/// Optional outer namespace(s)
 
44
CUB_NS_PREFIX
 
45
 
 
46
/// CUB namespace
 
47
namespace cub {
 
48
 
 
49
 
 
50
/**
 
51
 * \addtogroup WarpModule
 
52
 * @{
 
53
 */
 
54
 
 
55
/**
 
56
 * \brief The WarpReduce class provides [<em>collective</em>](index.html#sec0) methods for computing a parallel reduction of items partitioned across CUDA warp threads. ![](warp_reduce_logo.png)
 
57
 *
 
58
 * \par Overview
 
59
 * A <a href="http://en.wikipedia.org/wiki/Reduce_(higher-order_function)"><em>reduction</em></a> (or <em>fold</em>)
 
60
 * uses a binary combining operator to compute a single aggregate from a list of input elements.
 
61
 *
 
62
 * \tparam T                        The reduction input/output element type
 
63
 * \tparam LOGICAL_WARPS            <b>[optional]</b> The number of entrant "logical" warps performing concurrent warp reductions.  Default is 1.
 
64
 * \tparam LOGICAL_WARP_THREADS     <b>[optional]</b> The number of threads per "logical" warp (may be less than the number of hardware warp threads).  Default is the warp size of the targeted CUDA compute-capability (e.g., 32 threads for SM20).
 
65
 *
 
66
 * \par Simple Examples
 
67
 * \warpcollective{WarpReduce}
 
68
 * \par
 
69
 * The code snippet below illustrates four concurrent warp sum reductions within a block of
 
70
 * 128 threads (one per each of the 32-thread warps).
 
71
 * \par
 
72
 * \code
 
73
 * #include <cub/cub.cuh>
 
74
 *
 
75
 * __global__ void ExampleKernel(...)
 
76
 * {
 
77
 *     // Specialize WarpReduce for 4 warps on type int
 
78
 *     typedef cub::WarpReduce<int, 4> WarpReduce;
 
79
 *
 
80
 *     // Allocate shared memory for WarpReduce
 
81
 *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
82
 *
 
83
 *     // Obtain one input item per thread
 
84
 *     int thread_data = ...
 
85
 *
 
86
 *     // Return the warp-wide sums to each lane0 (threads 0, 32, 64, and 96)
 
87
 *     int aggregate = WarpReduce(temp_storage).Sum(thread_data);
 
88
 *
 
89
 * \endcode
 
90
 * \par
 
91
 * Suppose the set of input \p thread_data across the block of threads is <tt>0, 1, 2, 3, ..., 127</tt>.
 
92
 * The corresponding output \p aggregate in threads 0, 32, 64, and 96 will \p 496, \p 1520,
 
93
 * \p 2544, and \p 3568, respectively (and is undefined in other threads).
 
94
 *
 
95
 * \par
 
96
 * The code snippet below illustrates a single warp sum reduction within a block of
 
97
 * 128 threads.
 
98
 * \par
 
99
 * \code
 
100
 * #include <cub/cub.cuh>
 
101
 *
 
102
 * __global__ void ExampleKernel(...)
 
103
 * {
 
104
 *     // Specialize WarpReduce for one warp on type int
 
105
 *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
106
 *
 
107
 *     // Allocate shared memory for WarpReduce
 
108
 *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
109
 *     ...
 
110
 *
 
111
 *     // Only the first warp performs a reduction
 
112
 *     if (threadIdx.x < 32)
 
113
 *     {
 
114
 *         // Obtain one input item per thread
 
115
 *         int thread_data = ...
 
116
 *
 
117
 *         // Return the warp-wide sum to lane0
 
118
 *         int aggregate = WarpReduce(temp_storage).Sum(thread_data);
 
119
 *
 
120
 * \endcode
 
121
 * \par
 
122
 * Suppose the set of input \p thread_data across the warp of threads is <tt>0, 1, 2, 3, ..., 31</tt>.
 
123
 * The corresponding output \p aggregate in thread0 will be \p 496 (and is undefined in other threads).
 
124
 *
 
125
 * \par Usage and Performance Considerations
 
126
 * - Supports "logical" warps smaller than the physical warp size (e.g., logical warps of 8 threads)
 
127
 * - The number of entrant threads must be an multiple of \p LOGICAL_WARP_THREADS
 
128
 * - Warp reductions are concurrent if more than one logical warp is participating
 
129
 * - Uses special instructions when applicable (e.g., warp \p SHFL instructions)
 
130
 * - Uses synchronization-free communication between warp lanes when applicable
 
131
 * - Zero bank conflicts for most types
 
132
 * - Computation is slightly more efficient (i.e., having lower instruction overhead) for:
 
133
 *     - Summation (<b><em>vs.</em></b> generic reduction)
 
134
 *     - The architecture's warp size is a whole multiple of \p LOGICAL_WARP_THREADS
 
135
 *
 
136
 */
 
137
template <
 
138
    typename    T,
 
139
    int         LOGICAL_WARPS           = 1,
 
140
    int         LOGICAL_WARP_THREADS    = PtxArchProps::WARP_THREADS>
 
141
class WarpReduce
 
142
{
 
143
private:
 
144
 
 
145
    /******************************************************************************
 
146
     * Constants and typedefs
 
147
     ******************************************************************************/
 
148
 
 
149
    enum
 
150
    {
 
151
        POW_OF_TWO = ((LOGICAL_WARP_THREADS & (LOGICAL_WARP_THREADS - 1)) == 0),
 
152
    };
 
153
 
 
154
public:
 
155
 
 
156
    #ifndef DOXYGEN_SHOULD_SKIP_THIS    // Do not document
 
157
 
 
158
    /// Internal specialization.  Use SHFL-based reduction if (architecture is >= SM30) and ((only one logical warp) or (LOGICAL_WARP_THREADS is a power-of-two))
 
159
    typedef typename If<(CUB_PTX_ARCH >= 300) && ((LOGICAL_WARPS == 1) || POW_OF_TWO),
 
160
        WarpReduceShfl<T, LOGICAL_WARPS, LOGICAL_WARP_THREADS>,
 
161
        WarpReduceSmem<T, LOGICAL_WARPS, LOGICAL_WARP_THREADS> >::Type InternalWarpReduce;
 
162
 
 
163
    #endif // DOXYGEN_SHOULD_SKIP_THIS
 
164
 
 
165
 
 
166
private:
 
167
 
 
168
    /// Shared memory storage layout type for WarpReduce
 
169
    typedef typename InternalWarpReduce::TempStorage _TempStorage;
 
170
 
 
171
 
 
172
    /******************************************************************************
 
173
     * Thread fields
 
174
     ******************************************************************************/
 
175
 
 
176
    /// Shared storage reference
 
177
    _TempStorage &temp_storage;
 
178
 
 
179
    /// Warp ID
 
180
    int warp_id;
 
181
 
 
182
    /// Lane ID
 
183
    int lane_id;
 
184
 
 
185
 
 
186
    /******************************************************************************
 
187
     * Utility methods
 
188
     ******************************************************************************/
 
189
 
 
190
    /// Internal storage allocator
 
191
    __device__ __forceinline__ _TempStorage& PrivateStorage()
 
192
    {
 
193
        __shared__ TempStorage private_storage;
 
194
        return private_storage;
 
195
    }
 
196
 
 
197
 
 
198
public:
 
199
 
 
200
    /// \smemstorage{WarpReduce}
 
201
    struct TempStorage : Uninitialized<_TempStorage> {};
 
202
 
 
203
 
 
204
    /******************************************************************//**
 
205
     * \name Collective constructors
 
206
     *********************************************************************/
 
207
    //@{
 
208
 
 
209
 
 
210
    /**
 
211
     * \brief Collective constructor for 1D thread blocks using a private static allocation of shared memory as temporary storage.  Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
 
212
     *
 
213
     */
 
214
    __device__ __forceinline__ WarpReduce()
 
215
    :
 
216
        temp_storage(PrivateStorage()),
 
217
        warp_id((LOGICAL_WARPS == 1) ?
 
218
            0 :
 
219
            threadIdx.x / LOGICAL_WARP_THREADS),
 
220
        lane_id(((LOGICAL_WARPS == 1) || (LOGICAL_WARP_THREADS == PtxArchProps::WARP_THREADS)) ?
 
221
            LaneId() :
 
222
            threadIdx.x % LOGICAL_WARP_THREADS)
 
223
    {}
 
224
 
 
225
 
 
226
    /**
 
227
     * \brief Collective constructor for 1D thread blocks using the specified memory allocation as temporary storage.  Logical warp and lane identifiers are constructed from <tt>threadIdx.x</tt>.
 
228
     */
 
229
    __device__ __forceinline__ WarpReduce(
 
230
        TempStorage &temp_storage)             ///< [in] Reference to memory allocation having layout type TempStorage
 
231
    :
 
232
        temp_storage(temp_storage.Alias()),
 
233
        warp_id((LOGICAL_WARPS == 1) ?
 
234
            0 :
 
235
            threadIdx.x / LOGICAL_WARP_THREADS),
 
236
        lane_id(((LOGICAL_WARPS == 1) || (LOGICAL_WARP_THREADS == PtxArchProps::WARP_THREADS)) ?
 
237
            LaneId() :
 
238
            threadIdx.x % LOGICAL_WARP_THREADS)
 
239
    {}
 
240
 
 
241
 
 
242
    /**
 
243
     * \brief Collective constructor using a private static allocation of shared memory as temporary storage.  Threads are identified using the given warp and lane identifiers.
 
244
     */
 
245
    __device__ __forceinline__ WarpReduce(
 
246
        int warp_id,                           ///< [in] A suitable warp membership identifier
 
247
        int lane_id)                           ///< [in] A lane identifier within the warp
 
248
    :
 
249
        temp_storage(PrivateStorage()),
 
250
        warp_id(warp_id),
 
251
        lane_id(lane_id)
 
252
    {}
 
253
 
 
254
 
 
255
    /**
 
256
     * \brief Collective constructor using the specified memory allocation as temporary storage.  Threads are identified using the given warp and lane identifiers.
 
257
     */
 
258
    __device__ __forceinline__ WarpReduce(
 
259
        TempStorage &temp_storage,             ///< [in] Reference to memory allocation having layout type TempStorage
 
260
        int warp_id,                           ///< [in] A suitable warp membership identifier
 
261
        int lane_id)                           ///< [in] A lane identifier within the warp
 
262
    :
 
263
        temp_storage(temp_storage.Alias()),
 
264
        warp_id(warp_id),
 
265
        lane_id(lane_id)
 
266
    {}
 
267
 
 
268
 
 
269
 
 
270
    //@}  end member group
 
271
    /******************************************************************//**
 
272
     * \name Summation reductions
 
273
     *********************************************************************/
 
274
    //@{
 
275
 
 
276
 
 
277
    /**
 
278
     * \brief Computes a warp-wide sum in each active warp.  The output is valid in warp <em>lane</em><sub>0</sub>.
 
279
     *
 
280
     * \smemreuse
 
281
     *
 
282
     * The code snippet below illustrates four concurrent warp sum reductions within a block of
 
283
     * 128 threads (one per each of the 32-thread warps).
 
284
     * \par
 
285
     * \code
 
286
     * #include <cub/cub.cuh>
 
287
     *
 
288
     * __global__ void ExampleKernel(...)
 
289
     * {
 
290
     *     // Specialize WarpReduce for 4 warps on type int
 
291
     *     typedef cub::WarpReduce<int, 4> WarpReduce;
 
292
     *
 
293
     *     // Allocate shared memory for WarpReduce
 
294
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
295
     *
 
296
     *     // Obtain one input item per thread
 
297
     *     int thread_data = ...
 
298
     *
 
299
     *     // Return the warp-wide sums to each lane0
 
300
     *     int aggregate = WarpReduce(temp_storage).Sum(thread_data);
 
301
     *
 
302
     * \endcode
 
303
     * \par
 
304
     * Suppose the set of input \p thread_data across the block of threads is <tt>0, 1, 2, 3, ..., 127</tt>.
 
305
     * The corresponding output \p aggregate in threads 0, 32, 64, and 96 will \p 496, \p 1520,
 
306
     * \p 2544, and \p 3568, respectively (and is undefined in other threads).
 
307
     *
 
308
     */
 
309
    __device__ __forceinline__ T Sum(
 
310
        T                   input)              ///< [in] Calling thread's input
 
311
    {
 
312
        return InternalWarpReduce(temp_storage, warp_id, lane_id).Sum<true, 1>(input, LOGICAL_WARP_THREADS);
 
313
    }
 
314
 
 
315
    /**
 
316
     * \brief Computes a partially-full warp-wide sum in each active warp.  The output is valid in warp <em>lane</em><sub>0</sub>.
 
317
     *
 
318
     * All threads in each logical warp must agree on the same value for \p valid_items.  Otherwise the result is undefined.
 
319
     *
 
320
     * \smemreuse
 
321
     *
 
322
     * The code snippet below illustrates a sum reduction within a single, partially-full
 
323
     * block of 32 threads (one warp).
 
324
     * \par
 
325
     * \code
 
326
     * #include <cub/cub.cuh>
 
327
     *
 
328
     * __global__ void ExampleKernel(int *d_data, int valid_items)
 
329
     * {
 
330
     *     // Specialize WarpReduce for a single warp on type int
 
331
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
332
     *
 
333
     *     // Allocate shared memory for WarpReduce
 
334
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
335
     *
 
336
     *     // Obtain one input item per thread if in range
 
337
     *     int thread_data;
 
338
     *     if (threadIdx.x < valid_items)
 
339
     *         thread_data = d_data[threadIdx.x];
 
340
     *
 
341
     *     // Return the warp-wide sums to each lane0
 
342
     *     int aggregate = WarpReduce(temp_storage).Sum(
 
343
     *         thread_data, valid_items);
 
344
     *
 
345
     * \endcode
 
346
     * \par
 
347
     * Suppose the input \p d_data is <tt>0, 1, 2, 3, 4, ...</tt> and \p valid_items
 
348
     * is \p 4.  The corresponding output \p aggregate in thread0 is \p 6 (and is
 
349
     * undefined in other threads).
 
350
     *
 
351
     */
 
352
    __device__ __forceinline__ T Sum(
 
353
        T                   input,              ///< [in] Calling thread's input
 
354
        int                 valid_items)        ///< [in] Total number of valid items in the calling thread's logical warp (may be less than \p LOGICAL_WARP_THREADS)
 
355
    {
 
356
        // Determine if we don't need bounds checking
 
357
        if (valid_items >= LOGICAL_WARP_THREADS)
 
358
        {
 
359
            return InternalWarpReduce(temp_storage, warp_id, lane_id).Sum<true, 1>(input, valid_items);
 
360
        }
 
361
        else
 
362
        {
 
363
            return InternalWarpReduce(temp_storage, warp_id, lane_id).Sum<false, 1>(input, valid_items);
 
364
        }
 
365
    }
 
366
 
 
367
 
 
368
    /**
 
369
     * \brief Computes a segmented sum in each active warp where segments are defined by head-flags.  The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
 
370
     *
 
371
     * \smemreuse
 
372
     *
 
373
     * The code snippet below illustrates a head-segmented warp sum
 
374
     * reduction within a block of 32 threads (one warp).
 
375
     * \par
 
376
     * \code
 
377
     * #include <cub/cub.cuh>
 
378
     *
 
379
     * __global__ void ExampleKernel(...)
 
380
     * {
 
381
     *     // Specialize WarpReduce for a single warp on type int
 
382
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
383
     *
 
384
     *     // Allocate shared memory for WarpReduce
 
385
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
386
     *
 
387
     *     // Obtain one input item and flag per thread
 
388
     *     int thread_data = ...
 
389
     *     int head_flag = ...
 
390
     *
 
391
     *     // Return the warp-wide sums to each lane0
 
392
     *     int aggregate = WarpReduce(temp_storage).HeadSegmentedSum(
 
393
     *         thread_data, head_flag);
 
394
     *
 
395
     * \endcode
 
396
     * \par
 
397
     * Suppose the set of input \p thread_data and \p head_flag across the block of threads
 
398
     * is <tt>0, 1, 2, 3, ..., 31</tt> and is <tt>1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0</tt>,
 
399
     * respectively.  The corresponding output \p aggregate in threads 0, 4, 8, etc. will be
 
400
     * \p 6, \p 22, \p 38, etc. (and is undefined in other threads).
 
401
     *
 
402
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
403
     *
 
404
     */
 
405
    template <
 
406
        typename            Flag>
 
407
    __device__ __forceinline__ T HeadSegmentedSum(
 
408
        T                   input,              ///< [in] Calling thread's input
 
409
        Flag                head_flag)          ///< [in] Head flag denoting whether or not \p input is the start of a new segment
 
410
    {
 
411
        return HeadSegmentedReduce(input, head_flag, cub::Sum());
 
412
    }
 
413
 
 
414
 
 
415
    /**
 
416
     * \brief Computes a segmented sum in each active warp where segments are defined by tail-flags.  The sum of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
 
417
     *
 
418
     * \smemreuse
 
419
     *
 
420
     * The code snippet below illustrates a tail-segmented warp sum
 
421
     * reduction within a block of 32 threads (one warp).
 
422
     * \par
 
423
     * \code
 
424
     * #include <cub/cub.cuh>
 
425
     *
 
426
     * __global__ void ExampleKernel(...)
 
427
     * {
 
428
     *     // Specialize WarpReduce for a single warp on type int
 
429
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
430
     *
 
431
     *     // Allocate shared memory for WarpReduce
 
432
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
433
     *
 
434
     *     // Obtain one input item and flag per thread
 
435
     *     int thread_data = ...
 
436
     *     int tail_flag = ...
 
437
     *
 
438
     *     // Return the warp-wide sums to each lane0
 
439
     *     int aggregate = WarpReduce(temp_storage).TailSegmentedSum(
 
440
     *         thread_data, tail_flag);
 
441
     *
 
442
     * \endcode
 
443
     * \par
 
444
     * Suppose the set of input \p thread_data and \p tail_flag across the block of threads
 
445
     * is <tt>0, 1, 2, 3, ..., 31</tt> and is <tt>0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1</tt>,
 
446
     * respectively.  The corresponding output \p aggregate in threads 0, 4, 8, etc. will be
 
447
     * \p 6, \p 22, \p 38, etc. (and is undefined in other threads).
 
448
     *
 
449
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
450
     */
 
451
    template <
 
452
        typename            Flag>
 
453
    __device__ __forceinline__ T TailSegmentedSum(
 
454
        T                   input,              ///< [in] Calling thread's input
 
455
        Flag                tail_flag)          ///< [in] Head flag denoting whether or not \p input is the start of a new segment
 
456
    {
 
457
        return TailSegmentedReduce(input, tail_flag, cub::Sum());
 
458
    }
 
459
 
 
460
 
 
461
 
 
462
    //@}  end member group
 
463
    /******************************************************************//**
 
464
     * \name Generic reductions
 
465
     *********************************************************************/
 
466
    //@{
 
467
 
 
468
    /**
 
469
     * \brief Computes a warp-wide reduction in each active warp using the specified binary reduction functor.  The output is valid in warp <em>lane</em><sub>0</sub>.
 
470
     *
 
471
     * Supports non-commutative reduction operators
 
472
     *
 
473
     * \smemreuse
 
474
     *
 
475
     * The code snippet below illustrates four concurrent warp max reductions within a block of
 
476
     * 128 threads (one per each of the 32-thread warps).
 
477
     * \par
 
478
     * \code
 
479
     * #include <cub/cub.cuh>
 
480
     *
 
481
     * __global__ void ExampleKernel(...)
 
482
     * {
 
483
     *     // Specialize WarpReduce for 4 warps on type int
 
484
     *     typedef cub::WarpReduce<int, 4> WarpReduce;
 
485
     *
 
486
     *     // Allocate shared memory for WarpReduce
 
487
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
488
     *
 
489
     *     // Obtain one input item per thread
 
490
     *     int thread_data = ...
 
491
     *
 
492
     *     // Return the warp-wide reductions to each lane0
 
493
     *     int aggregate = WarpReduce(temp_storage).Reduce(
 
494
     *         thread_data, cub::Max());
 
495
     *
 
496
     * \endcode
 
497
     * \par
 
498
     * Suppose the set of input \p thread_data across the block of threads is <tt>0, 1, 2, 3, ..., 127</tt>.
 
499
     * The corresponding output \p aggregate in threads 0, 32, 64, and 96 will \p 31, \p 63,
 
500
     * \p 95, and \p 127, respectively  (and is undefined in other threads).
 
501
     *
 
502
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
503
     */
 
504
    template <typename ReductionOp>
 
505
    __device__ __forceinline__ T Reduce(
 
506
        T                   input,              ///< [in] Calling thread's input
 
507
        ReductionOp         reduction_op)       ///< [in] Binary reduction operator
 
508
    {
 
509
        return InternalWarpReduce(temp_storage, warp_id, lane_id).Reduce<true, 1>(input, LOGICAL_WARP_THREADS, reduction_op);
 
510
    }
 
511
 
 
512
    /**
 
513
     * \brief Computes a partially-full warp-wide reduction in each active warp using the specified binary reduction functor.  The output is valid in warp <em>lane</em><sub>0</sub>.
 
514
     *
 
515
     * All threads in each logical warp must agree on the same value for \p valid_items.  Otherwise the result is undefined.
 
516
     *
 
517
     * Supports non-commutative reduction operators
 
518
     *
 
519
     * \smemreuse
 
520
     *
 
521
     * The code snippet below illustrates a max reduction within a single, partially-full
 
522
     * block of 32 threads (one warp).
 
523
     * \par
 
524
     * \code
 
525
     * #include <cub/cub.cuh>
 
526
     *
 
527
     * __global__ void ExampleKernel(int *d_data, int valid_items)
 
528
     * {
 
529
     *     // Specialize WarpReduce for a single warp on type int
 
530
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
531
     *
 
532
     *     // Allocate shared memory for WarpReduce
 
533
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
534
     *
 
535
     *     // Obtain one input item per thread if in range
 
536
     *     int thread_data;
 
537
     *     if (threadIdx.x < valid_items)
 
538
     *         thread_data = d_data[threadIdx.x];
 
539
     *
 
540
     *     // Return the warp-wide reductions to each lane0
 
541
     *     int aggregate = WarpReduce(temp_storage).Reduce(
 
542
     *         thread_data, cub::Max(), valid_items);
 
543
     *
 
544
     * \endcode
 
545
     * \par
 
546
     * Suppose the input \p d_data is <tt>0, 1, 2, 3, 4, ...</tt> and \p valid_items
 
547
     * is \p 4.  The corresponding output \p aggregate in thread0 is \p 3 (and is
 
548
     * undefined in other threads).
 
549
     *
 
550
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
551
     */
 
552
    template <typename ReductionOp>
 
553
    __device__ __forceinline__ T Reduce(
 
554
        T                   input,              ///< [in] Calling thread's input
 
555
        ReductionOp         reduction_op,       ///< [in] Binary reduction operator
 
556
        int                 valid_items)        ///< [in] Total number of valid items in the calling thread's logical warp (may be less than \p LOGICAL_WARP_THREADS)
 
557
    {
 
558
        // Determine if we don't need bounds checking
 
559
        if (valid_items >= LOGICAL_WARP_THREADS)
 
560
        {
 
561
            return InternalWarpReduce(temp_storage, warp_id, lane_id).Reduce<true, 1>(input, valid_items, reduction_op);
 
562
        }
 
563
        else
 
564
        {
 
565
            return InternalWarpReduce(temp_storage, warp_id, lane_id).Reduce<false, 1>(input, valid_items, reduction_op);
 
566
        }
 
567
    }
 
568
 
 
569
 
 
570
    /**
 
571
     * \brief Computes a segmented reduction in each active warp where segments are defined by head-flags.  The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
 
572
     *
 
573
     * Supports non-commutative reduction operators
 
574
     *
 
575
     * \smemreuse
 
576
     *
 
577
     * The code snippet below illustrates a head-segmented warp max
 
578
     * reduction within a block of 32 threads (one warp).
 
579
     * \par
 
580
     * \code
 
581
     * #include <cub/cub.cuh>
 
582
     *
 
583
     * __global__ void ExampleKernel(...)
 
584
     * {
 
585
     *     // Specialize WarpReduce for a single warp on type int
 
586
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
587
     *
 
588
     *     // Allocate shared memory for WarpReduce
 
589
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
590
     *
 
591
     *     // Obtain one input item and flag per thread
 
592
     *     int thread_data = ...
 
593
     *     int head_flag = ...
 
594
     *
 
595
     *     // Return the warp-wide reductions to each lane0
 
596
     *     int aggregate = WarpReduce(temp_storage).HeadSegmentedReduce(
 
597
     *         thread_data, head_flag, cub::Max());
 
598
     *
 
599
     * \endcode
 
600
     * \par
 
601
     * Suppose the set of input \p thread_data and \p head_flag across the block of threads
 
602
     * is <tt>0, 1, 2, 3, ..., 31</tt> and is <tt>1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0</tt>,
 
603
     * respectively.  The corresponding output \p aggregate in threads 0, 4, 8, etc. will be
 
604
     * \p 3, \p 7, \p 11, etc. (and is undefined in other threads).
 
605
     *
 
606
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
607
     */
 
608
    template <
 
609
        typename            ReductionOp,
 
610
        typename            Flag>
 
611
    __device__ __forceinline__ T HeadSegmentedReduce(
 
612
        T                   input,              ///< [in] Calling thread's input
 
613
        Flag                head_flag,          ///< [in] Head flag denoting whether or not \p input is the start of a new segment
 
614
        ReductionOp         reduction_op)       ///< [in] Reduction operator
 
615
    {
 
616
        return InternalWarpReduce(temp_storage, warp_id, lane_id).template SegmentedReduce<true>(input, head_flag, reduction_op);
 
617
    }
 
618
 
 
619
 
 
620
    /**
 
621
     * \brief Computes a segmented reduction in each active warp where segments are defined by tail-flags.  The reduction of each segment is returned to the first lane in that segment (which always includes <em>lane</em><sub>0</sub>).
 
622
     *
 
623
     * Supports non-commutative reduction operators
 
624
     *
 
625
     * \smemreuse
 
626
     *
 
627
     * The code snippet below illustrates a tail-segmented warp max
 
628
     * reduction within a block of 32 threads (one warp).
 
629
     * \par
 
630
     * \code
 
631
     * #include <cub/cub.cuh>
 
632
     *
 
633
     * __global__ void ExampleKernel(...)
 
634
     * {
 
635
     *     // Specialize WarpReduce for a single warp on type int
 
636
     *     typedef cub::WarpReduce<int, 1> WarpReduce;
 
637
     *
 
638
     *     // Allocate shared memory for WarpReduce
 
639
     *     __shared__ typename WarpReduce::TempStorage temp_storage;
 
640
     *
 
641
     *     // Obtain one input item and flag per thread
 
642
     *     int thread_data = ...
 
643
     *     int tail_flag = ...
 
644
     *
 
645
     *     // Return the warp-wide reductions to each lane0
 
646
     *     int aggregate = WarpReduce(temp_storage).TailSegmentedReduce(
 
647
     *         thread_data, tail_flag, cub::Max());
 
648
     *
 
649
     * \endcode
 
650
     * \par
 
651
     * Suppose the set of input \p thread_data and \p tail_flag across the block of threads
 
652
     * is <tt>0, 1, 2, 3, ..., 31</tt> and is <tt>0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1</tt>,
 
653
     * respectively.  The corresponding output \p aggregate in threads 0, 4, 8, etc. will be
 
654
     * \p 3, \p 7, \p 11, etc. (and is undefined in other threads).
 
655
     *
 
656
     * \tparam ReductionOp     <b>[inferred]</b> Binary reduction operator type having member <tt>T operator()(const T &a, const T &b)</tt>
 
657
     */
 
658
    template <
 
659
        typename            ReductionOp,
 
660
        typename            Flag>
 
661
    __device__ __forceinline__ T TailSegmentedReduce(
 
662
        T                   input,              ///< [in] Calling thread's input
 
663
        Flag                tail_flag,          ///< [in] Tail flag denoting whether or not \p input is the end of the current segment
 
664
        ReductionOp         reduction_op)       ///< [in] Reduction operator
 
665
    {
 
666
        return InternalWarpReduce(temp_storage, warp_id, lane_id).template SegmentedReduce<false>(input, tail_flag, reduction_op);
 
667
    }
 
668
 
 
669
 
 
670
 
 
671
    //@}  end member group
 
672
};
 
673
 
 
674
/** @} */       // end group WarpModule
 
675
 
 
676
}               // CUB namespace
 
677
CUB_NS_POSTFIX  // Optional outer namespace(s)