← Back to branch summary

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

« back to all changes in this revision

Viewing changes to lib/kokkos/TPL/cub/device/block/specializations/block_histo_tiles_satomic.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/device/block/specializations/block_histo_tiles_satomic.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::BlockHistogramTilesSharedAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using shared atomics
 
32
 */
 
33
 
 
34
#pragma once
 
35
 
 
36
#include <iterator>
 
37
 
 
38
#include "../../../util_type.cuh"
 
39
#include "../../../util_namespace.cuh"
 
40
 
 
41
/// Optional outer namespace(s)
 
42
CUB_NS_PREFIX
 
43
 
 
44
/// CUB namespace
 
45
namespace cub {
 
46
 
 
47
 
 
48
/**
 
49
 * BlockHistogramTilesSharedAtomic implements a stateful abstraction of CUDA thread blocks for histogramming multiple tiles as part of device-wide histogram using shared atomics
 
50
 */
 
51
template <
 
52
    typename    BlockHistogramTilesPolicy,          ///< Tuning policy
 
53
    int         BINS,                           ///< Number of histogram bins
 
54
    int         CHANNELS,                       ///< Number of channels interleaved in the input data (may be greater than the number of active channels being histogrammed)
 
55
    int         ACTIVE_CHANNELS,                ///< Number of channels actively being histogrammed
 
56
    typename    InputIteratorRA,                ///< The input iterator type (may be a simple pointer type).  Must have a value type that can be cast as an integer in the range [0..BINS-1]
 
57
    typename    HistoCounter,                   ///< Integral type for counting sample occurrences per histogram bin
 
58
    typename    SizeT>                          ///< Integer type for offsets
 
59
struct BlockHistogramTilesSharedAtomic
 
60
{
 
61
    //---------------------------------------------------------------------
 
62
    // Types and constants
 
63
    //---------------------------------------------------------------------
 
64
 
 
65
    // Sample type
 
66
    typedef typename std::iterator_traits<InputIteratorRA>::value_type SampleT;
 
67
 
 
68
    // Constants
 
69
    enum
 
70
    {
 
71
        BLOCK_THREADS       = BlockHistogramTilesPolicy::BLOCK_THREADS,
 
72
        ITEMS_PER_THREAD    = BlockHistogramTilesPolicy::ITEMS_PER_THREAD,
 
73
        TILE_CHANNEL_ITEMS  = BLOCK_THREADS * ITEMS_PER_THREAD,
 
74
        TILE_ITEMS          = TILE_CHANNEL_ITEMS * CHANNELS,
 
75
    };
 
76
 
 
77
    /// Shared memory type required by this thread block
 
78
    struct _TempStorage
 
79
    {
 
80
        HistoCounter histograms[ACTIVE_CHANNELS][BINS + 1];  // One word of padding between channel histograms to prevent warps working on different histograms from hammering on the same bank
 
81
    };
 
82
 
 
83
 
 
84
    /// Alias wrapper allowing storage to be unioned
 
85
    struct TempStorage : Uninitialized<_TempStorage> {};
 
86
 
 
87
 
 
88
    //---------------------------------------------------------------------
 
89
    // Per-thread fields
 
90
    //---------------------------------------------------------------------
 
91
 
 
92
    /// Reference to temp_storage
 
93
    _TempStorage &temp_storage;
 
94
 
 
95
    /// Reference to output histograms
 
96
    HistoCounter* (&d_out_histograms)[ACTIVE_CHANNELS];
 
97
 
 
98
    /// Input data to reduce
 
99
    InputIteratorRA d_in;
 
100
 
 
101
 
 
102
    //---------------------------------------------------------------------
 
103
    // Interface
 
104
    //---------------------------------------------------------------------
 
105
 
 
106
    /**
 
107
     * Constructor
 
108
     */
 
109
    __device__ __forceinline__ BlockHistogramTilesSharedAtomic(
 
110
        TempStorage         &temp_storage,                                  ///< Reference to temp_storage
 
111
        InputIteratorRA     d_in,                                           ///< Input data to reduce
 
112
        HistoCounter*       (&d_out_histograms)[ACTIVE_CHANNELS])           ///< Reference to output histograms
 
113
    :
 
114
        temp_storage(temp_storage.Alias()),
 
115
        d_in(d_in),
 
116
        d_out_histograms(d_out_histograms)
 
117
    {
 
118
        // Initialize histogram bin counts to zeros
 
119
        #pragma unroll
 
120
        for (int CHANNEL = 0; CHANNEL < ACTIVE_CHANNELS; ++CHANNEL)
 
121
        {
 
122
            int histo_offset = 0;
 
123
 
 
124
            #pragma unroll
 
125
            for(; histo_offset + BLOCK_THREADS <= BINS; histo_offset += BLOCK_THREADS)
 
126
            {
 
127
                this->temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x] = 0;
 
128
            }
 
129
            // Finish up with guarded initialization if necessary
 
130
            if ((BINS % BLOCK_THREADS != 0) && (histo_offset + threadIdx.x < BINS))
 
131
            {
 
132
                this->temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x] = 0;
 
133
            }
 
134
        }
 
135
    }
 
136
 
 
137
 
 
138
    /**
 
139
     * Process a single tile of input
 
140
     */
 
141
    template <bool FULL_TILE>
 
142
    __device__ __forceinline__ void ConsumeTile(
 
143
        SizeT   block_offset,               ///< The offset the tile to consume
 
144
        int     valid_items = TILE_ITEMS)   ///< The number of valid items in the tile
 
145
    {
 
146
        if (FULL_TILE)
 
147
        {
 
148
            // Full tile of samples to read and composite
 
149
            SampleT items[ITEMS_PER_THREAD][CHANNELS];
 
150
 
 
151
            #pragma unroll
 
152
            for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
 
153
            {
 
154
                #pragma unroll
 
155
                for (int CHANNEL = 0; CHANNEL < CHANNELS; ++CHANNEL)
 
156
                {
 
157
                    if (CHANNEL < ACTIVE_CHANNELS)
 
158
                    {
 
159
                        items[ITEM][CHANNEL] = d_in[block_offset + (ITEM * BLOCK_THREADS * CHANNELS) + (threadIdx.x * CHANNELS) + CHANNEL];
 
160
                    }
 
161
                }
 
162
            }
 
163
 
 
164
            __threadfence_block();
 
165
 
 
166
            #pragma unroll
 
167
            for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ITEM++)
 
168
            {
 
169
                #pragma unroll
 
170
                for (int CHANNEL = 0; CHANNEL < CHANNELS; ++CHANNEL)
 
171
                {
 
172
                    if (CHANNEL < ACTIVE_CHANNELS)
 
173
                    {
 
174
                        atomicAdd(temp_storage.histograms[CHANNEL] + items[ITEM][CHANNEL], 1);
 
175
                    }
 
176
                }
 
177
            }
 
178
 
 
179
            __threadfence_block();
 
180
        }
 
181
        else
 
182
        {
 
183
            // Only a partially-full tile of samples to read and composite
 
184
            int bounds = valid_items - (threadIdx.x * CHANNELS);
 
185
 
 
186
            #pragma unroll
 
187
            for (int ITEM = 0; ITEM < ITEMS_PER_THREAD; ++ITEM)
 
188
            {
 
189
                #pragma unroll
 
190
                for (int CHANNEL = 0; CHANNEL < CHANNELS; ++CHANNEL)
 
191
                {
 
192
                    if (((ACTIVE_CHANNELS == CHANNELS) || (CHANNEL < ACTIVE_CHANNELS)) && ((ITEM * BLOCK_THREADS * CHANNELS) + CHANNEL < bounds))
 
193
                    {
 
194
                        SampleT item  = d_in[block_offset + (ITEM * BLOCK_THREADS * CHANNELS) + (threadIdx.x * CHANNELS) + CHANNEL];
 
195
                        atomicAdd(temp_storage.histograms[CHANNEL] + item, 1);
 
196
                    }
 
197
                }
 
198
            }
 
199
 
 
200
        }
 
201
    }
 
202
 
 
203
 
 
204
    /**
 
205
     * Aggregate results into output
 
206
     */
 
207
    __device__ __forceinline__ void AggregateOutput()
 
208
    {
 
209
        // Barrier to ensure shared memory histograms are coherent
 
210
        __syncthreads();
 
211
 
 
212
        // Copy shared memory histograms to output
 
213
        #pragma unroll
 
214
        for (int CHANNEL = 0; CHANNEL < ACTIVE_CHANNELS; ++CHANNEL)
 
215
        {
 
216
            int channel_offset  = (blockIdx.x * BINS);
 
217
            int histo_offset    = 0;
 
218
 
 
219
            #pragma unroll
 
220
            for(; histo_offset + BLOCK_THREADS <= BINS; histo_offset += BLOCK_THREADS)
 
221
            {
 
222
                d_out_histograms[CHANNEL][channel_offset + histo_offset + threadIdx.x] = temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x];
 
223
            }
 
224
            // Finish up with guarded initialization if necessary
 
225
            if ((BINS % BLOCK_THREADS != 0) && (histo_offset + threadIdx.x < BINS))
 
226
            {
 
227
                d_out_histograms[CHANNEL][channel_offset + histo_offset + threadIdx.x] = temp_storage.histograms[CHANNEL][histo_offset + threadIdx.x];
 
228
            }
 
229
        }
 
230
    }
 
231
};
 
232
 
 
233
 
 
234
 
 
235
}               // CUB namespace
 
236
CUB_NS_POSTFIX  // Optional outer namespace(s)
 
237