← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/kokkos/TPL/cub/block/specializations/block_scan_warp_scans.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/block/specializations/block_scan_warp_scans.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
 * cub::BlockScanWarpscans provides warpscan-based variants of parallel prefix scan across a CUDA threadblock.
 
32
 */
 
33
 
 
34
#pragma once
 
35
 
 
36
#include "../../util_arch.cuh"
 
37
#include "../../warp/warp_scan.cuh"
 
38
#include "../../util_namespace.cuh"
 
39
 
 
40
/// Optional outer namespace(s)
 
41
CUB_NS_PREFIX
 
42
 
 
43
/// CUB namespace
 
44
namespace cub {
 
45
 
 
46
/**
 
47
 * \brief BlockScanWarpScans provides warpscan-based variants of parallel prefix scan across a CUDA threadblock.
 
48
 */
 
49
template <
 
50
    typename            T,
 
51
    int                 BLOCK_THREADS>
 
52
struct BlockScanWarpScans
 
53
{
 
54
    /// Constants
 
55
    enum
 
56
    {
 
57
        /// Number of active warps
 
58
        WARPS = (BLOCK_THREADS + PtxArchProps::WARP_THREADS - 1) / PtxArchProps::WARP_THREADS,
 
59
    };
 
60
 
 
61
    ///  WarpScan utility type
 
62
    typedef WarpScan<T, WARPS, PtxArchProps::WARP_THREADS> WarpScan;
 
63
 
 
64
    /// Shared memory storage layout type
 
65
    struct _TempStorage
 
66
    {
 
67
        typename WarpScan::TempStorage      warp_scan;                  ///< Buffer for warp-synchronous scan
 
68
        T                                   warp_aggregates[WARPS];     ///< Shared totals from each warp-synchronous scan
 
69
        T                                   block_prefix;               ///< Shared prefix for the entire threadblock
 
70
    };
 
71
 
 
72
 
 
73
    /// Alias wrapper allowing storage to be unioned
 
74
    struct TempStorage : Uninitialized<_TempStorage> {};
 
75
 
 
76
 
 
77
    // Thread fields
 
78
    _TempStorage &temp_storage;
 
79
    int linear_tid;
 
80
    int warp_id;
 
81
    int lane_id;
 
82
 
 
83
 
 
84
    /// Constructor
 
85
    __device__ __forceinline__ BlockScanWarpScans(
 
86
        TempStorage &temp_storage,
 
87
        int linear_tid)
 
88
    :
 
89
        temp_storage(temp_storage.Alias()),
 
90
        linear_tid(linear_tid),
 
91
        warp_id((BLOCK_THREADS <= PtxArchProps::WARP_THREADS) ?
 
92
            0 :
 
93
            linear_tid / PtxArchProps::WARP_THREADS),
 
94
        lane_id((BLOCK_THREADS <= PtxArchProps::WARP_THREADS) ?
 
95
            linear_tid :
 
96
            linear_tid % PtxArchProps::WARP_THREADS)
 
97
    {}
 
98
 
 
99
 
 
100
    /// Update the calling thread's partial reduction with the warp-wide aggregates from preceding warps.  Also returns block-wide aggregate in <em>thread</em><sub>0</sub>.
 
101
    template <typename ScanOp>
 
102
    __device__ __forceinline__ void ApplyWarpAggregates(
 
103
        T               &partial,           ///< [out] The calling thread's partial reduction
 
104
        ScanOp          scan_op,            ///< [in] Binary scan operator
 
105
        T               warp_aggregate,     ///< [in] <b>[<em>lane</em><sub>0</sub>s only]</b> Warp-wide aggregate reduction of input items
 
106
        T               &block_aggregate,   ///< [out] Threadblock-wide aggregate reduction of input items
 
107
        bool            lane_valid = true)  ///< [in] Whether or not the partial belonging to the current thread is valid
 
108
    {
 
109
        // Share lane aggregates
 
110
        temp_storage.warp_aggregates[warp_id] = warp_aggregate;
 
111
 
 
112
        __syncthreads();
 
113
 
 
114
        block_aggregate = temp_storage.warp_aggregates[0];
 
115
 
 
116
        #pragma unroll
 
117
        for (int WARP = 1; WARP < WARPS; WARP++)
 
118
        {
 
119
            if (warp_id == WARP)
 
120
            {
 
121
                partial = (lane_valid) ?
 
122
                    scan_op(block_aggregate, partial) :     // fold it in our valid partial
 
123
                    block_aggregate;                        // replace our invalid partial with the aggregate
 
124
            }
 
125
 
 
126
            block_aggregate = scan_op(block_aggregate, temp_storage.warp_aggregates[WARP]);
 
127
        }
 
128
    }
 
129
 
 
130
 
 
131
    /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
132
    template <typename ScanOp>
 
133
    __device__ __forceinline__ void ExclusiveScan(
 
134
        T               input,              ///< [in] Calling thread's input items
 
135
        T               &output,            ///< [out] Calling thread's output items (may be aliased to \p input)
 
136
        const T         &identity,          ///< [in] Identity value
 
137
        ScanOp          scan_op,            ///< [in] Binary scan operator
 
138
        T               &block_aggregate)   ///< [out] Threadblock-wide aggregate reduction of input items
 
139
    {
 
140
        T warp_aggregate;
 
141
        WarpScan(temp_storage.warp_scan, warp_id, lane_id).ExclusiveScan(input, output, identity, scan_op, warp_aggregate);
 
142
 
 
143
        // Update outputs and block_aggregate with warp-wide aggregates
 
144
        ApplyWarpAggregates(output, scan_op, warp_aggregate, block_aggregate);
 
145
    }
 
146
 
 
147
 
 
148
    /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
149
    template <
 
150
        typename ScanOp,
 
151
        typename BlockPrefixOp>
 
152
    __device__ __forceinline__ void ExclusiveScan(
 
153
        T               input,                          ///< [in] Calling thread's input item
 
154
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
155
        T               identity,                       ///< [in] Identity value
 
156
        ScanOp          scan_op,                        ///< [in] Binary scan operator
 
157
        T               &block_aggregate,               ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
 
158
        BlockPrefixOp   &block_prefix_op)               ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
 
159
    {
 
160
        ExclusiveScan(input, output, identity, scan_op, block_aggregate);
 
161
 
 
162
        // Compute and share threadblock prefix
 
163
        if (warp_id == 0)
 
164
        {
 
165
            temp_storage.block_prefix = block_prefix_op(block_aggregate);
 
166
        }
 
167
 
 
168
        __syncthreads();
 
169
 
 
170
        // Incorporate threadblock prefix into outputs
 
171
        output = scan_op(temp_storage.block_prefix, output);
 
172
    }
 
173
 
 
174
 
 
175
    /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  Also provides every thread with the block-wide \p block_aggregate of all inputs.  With no identity value, the output computed for <em>thread</em><sub>0</sub> is undefined.
 
176
    template <typename ScanOp>
 
177
    __device__ __forceinline__ void ExclusiveScan(
 
178
        T               input,                          ///< [in] Calling thread's input item
 
179
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
180
        ScanOp          scan_op,                        ///< [in] Binary scan operator
 
181
        T               &block_aggregate)               ///< [out] Threadblock-wide aggregate reduction of input items
 
182
    {
 
183
        T warp_aggregate;
 
184
        WarpScan(temp_storage.warp_scan, warp_id, lane_id).ExclusiveScan(input, output, scan_op, warp_aggregate);
 
185
 
 
186
        // Update outputs and block_aggregate with warp-wide aggregates
 
187
        ApplyWarpAggregates(output, scan_op, warp_aggregate, block_aggregate, (lane_id > 0));
 
188
    }
 
189
 
 
190
 
 
191
    /// Computes an exclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
192
    template <
 
193
        typename ScanOp,
 
194
        typename BlockPrefixOp>
 
195
    __device__ __forceinline__ void ExclusiveScan(
 
196
        T               input,                          ///< [in] Calling thread's input item
 
197
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
198
        ScanOp          scan_op,                        ///< [in] Binary scan operator
 
199
        T               &block_aggregate,               ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
 
200
        BlockPrefixOp   &block_prefix_op)               ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
 
201
    {
 
202
        ExclusiveScan(input, output, scan_op, block_aggregate);
 
203
 
 
204
        // Compute and share threadblock prefix
 
205
        if (warp_id == 0)
 
206
        {
 
207
            temp_storage.block_prefix = block_prefix_op(block_aggregate);
 
208
        }
 
209
 
 
210
        __syncthreads();
 
211
 
 
212
        // Incorporate threadblock prefix into outputs
 
213
        output = (linear_tid == 0) ?
 
214
            temp_storage.block_prefix :
 
215
            scan_op(temp_storage.block_prefix, output);
 
216
    }
 
217
 
 
218
 
 
219
    /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator.  Each thread contributes one input element.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
220
    __device__ __forceinline__ void ExclusiveSum(
 
221
        T               input,                          ///< [in] Calling thread's input item
 
222
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
223
        T               &block_aggregate)               ///< [out] Threadblock-wide aggregate reduction of input items
 
224
    {
 
225
        T warp_aggregate;
 
226
        WarpScan(temp_storage.warp_scan, warp_id, lane_id).ExclusiveSum(input, output, warp_aggregate);
 
227
 
 
228
        // Update outputs and block_aggregate with warp-wide aggregates from lane-0s
 
229
        ApplyWarpAggregates(output, Sum(), warp_aggregate, block_aggregate);
 
230
    }
 
231
 
 
232
 
 
233
    /// Computes an exclusive threadblock-wide prefix scan using addition (+) as the scan operator.  Each thread contributes one input element.  Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
234
    template <typename BlockPrefixOp>
 
235
    __device__ __forceinline__ void ExclusiveSum(
 
236
        T               input,                          ///< [in] Calling thread's input item
 
237
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
238
        T               &block_aggregate,               ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
 
239
        BlockPrefixOp   &block_prefix_op)               ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
 
240
    {
 
241
        ExclusiveSum(input, output, block_aggregate);
 
242
 
 
243
        // Compute and share threadblock prefix
 
244
        if (warp_id == 0)
 
245
        {
 
246
            temp_storage.block_prefix = block_prefix_op(block_aggregate);
 
247
        }
 
248
 
 
249
        __syncthreads();
 
250
 
 
251
        // Incorporate threadblock prefix into outputs
 
252
        Sum scan_op;
 
253
        output = scan_op(temp_storage.block_prefix, output);
 
254
    }
 
255
 
 
256
 
 
257
    /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
258
    template <typename ScanOp>
 
259
    __device__ __forceinline__ void InclusiveScan(
 
260
        T               input,                          ///< [in] Calling thread's input item
 
261
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
262
        ScanOp          scan_op,                        ///< [in] Binary scan operator
 
263
        T               &block_aggregate)               ///< [out] Threadblock-wide aggregate reduction of input items
 
264
    {
 
265
        T warp_aggregate;
 
266
        WarpScan(temp_storage.warp_scan, warp_id, lane_id).InclusiveScan(input, output, scan_op, warp_aggregate);
 
267
 
 
268
        // Update outputs and block_aggregate with warp-wide aggregates from lane-0s
 
269
        ApplyWarpAggregates(output, scan_op, warp_aggregate, block_aggregate);
 
270
 
 
271
    }
 
272
 
 
273
 
 
274
    /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
275
    template <
 
276
        typename ScanOp,
 
277
        typename BlockPrefixOp>
 
278
    __device__ __forceinline__ void InclusiveScan(
 
279
        T               input,                          ///< [in] Calling thread's input item
 
280
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
281
        ScanOp          scan_op,                        ///< [in] Binary scan operator
 
282
        T               &block_aggregate,               ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
 
283
        BlockPrefixOp   &block_prefix_op)               ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
 
284
    {
 
285
        InclusiveScan(input, output, scan_op, block_aggregate);
 
286
 
 
287
        // Compute and share threadblock prefix
 
288
        if (warp_id == 0)
 
289
        {
 
290
            temp_storage.block_prefix = block_prefix_op(block_aggregate);
 
291
        }
 
292
 
 
293
        __syncthreads();
 
294
 
 
295
        // Incorporate threadblock prefix into outputs
 
296
        output = scan_op(temp_storage.block_prefix, output);
 
297
    }
 
298
 
 
299
 
 
300
    /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
301
    __device__ __forceinline__ void InclusiveSum(
 
302
        T               input,                          ///< [in] Calling thread's input item
 
303
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
304
        T               &block_aggregate)               ///< [out] Threadblock-wide aggregate reduction of input items
 
305
    {
 
306
        T warp_aggregate;
 
307
        WarpScan(temp_storage.warp_scan, warp_id, lane_id).InclusiveSum(input, output, warp_aggregate);
 
308
 
 
309
        // Update outputs and block_aggregate with warp-wide aggregates from lane-0s
 
310
        ApplyWarpAggregates(output, Sum(), warp_aggregate, block_aggregate);
 
311
    }
 
312
 
 
313
 
 
314
    /// Computes an inclusive threadblock-wide prefix scan using the specified binary \p scan_op functor.  Each thread contributes one input element.  Instead of using 0 as the threadblock-wide prefix, the call-back functor \p block_prefix_op is invoked by the first warp in the block, and the value returned by <em>lane</em><sub>0</sub> in that warp is used as the "seed" value that logically prefixes the threadblock's scan inputs.  Also provides every thread with the block-wide \p block_aggregate of all inputs.
 
315
    template <typename BlockPrefixOp>
 
316
    __device__ __forceinline__ void InclusiveSum(
 
317
        T               input,                          ///< [in] Calling thread's input item
 
318
        T               &output,                        ///< [out] Calling thread's output item (may be aliased to \p input)
 
319
        T               &block_aggregate,               ///< [out] Threadblock-wide aggregate reduction of input items (exclusive of the \p block_prefix_op value)
 
320
        BlockPrefixOp   &block_prefix_op)               ///< [in-out] <b>[<em>warp</em><sub>0</sub> only]</b> Call-back functor for specifying a threadblock-wide prefix to be applied to all inputs.
 
321
    {
 
322
        InclusiveSum(input, output, block_aggregate);
 
323
 
 
324
        // Compute and share threadblock prefix
 
325
        if (warp_id == 0)
 
326
        {
 
327
            temp_storage.block_prefix = block_prefix_op(block_aggregate);
 
328
        }
 
329
 
 
330
        __syncthreads();
 
331
 
 
332
        // Incorporate threadblock prefix into outputs
 
333
        Sum scan_op;
 
334
        output = scan_op(temp_storage.block_prefix, output);
 
335
    }
 
336
 
 
337
};
 
338
 
 
339
 
 
340
}               // CUB namespace
 
341
CUB_NS_POSTFIX  // Optional outer namespace(s)
 
342