/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.1/hipcub/include/hipcub/backend/rocprim/block/block_histogram.hpp Source File

• home
• docs
• checkouts
• readthedocs.org
• user_builds
• advanced-micro-devices-hipcub
• checkouts
• docs-5.5.1
• hipcub
• include
• hipcub
• backend
• rocprim
• block

block_histogram.hpp

1 /******************************************************************************

2  * Copyright (c) 2010-2011, Duane Merrill. All rights reserved.

3  * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.

4  * Modifications Copyright (c) 2017-2022, Advanced Micro Devices, Inc. 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 

30 #ifndef HIPCUB_ROCPRIM_BLOCK_BLOCK_HISTOGRAM_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_HISTOGRAM_HPP_

32 

33 #include "../util_ptx.hpp"

34 

35 #include <type_traits>

36 

37 #include <rocprim/block/block_histogram.hpp>

38 

39 BEGIN_HIPCUB_NAMESPACE

40 

41 namespace detail

42 {

43  inline constexpr

44  typename std::underlying_type<::rocprim::block_histogram_algorithm>::type

45  to_BlockHistogramAlgorithm_enum(::rocprim::block_histogram_algorithm v)

46  {

47  using utype = std::underlying_type<::rocprim::block_histogram_algorithm>::type;

48  return static_cast<utype>(v);

49  }

50 }

51 

52 enum BlockHistogramAlgorithm

53 {

54  BLOCK_HISTO_ATOMIC

55  = detail::to_BlockHistogramAlgorithm_enum(::rocprim::block_histogram_algorithm::using_atomic),

56  BLOCK_HISTO_SORT

57  = detail::to_BlockHistogramAlgorithm_enum(::rocprim::block_histogram_algorithm::using_sort)

58 };

59 

60 template<

61  typename T,

62  int BLOCK_DIM_X,

63  int ITEMS_PER_THREAD,

64  int BINS,

65  BlockHistogramAlgorithm ALGORITHM = BLOCK_HISTO_SORT,

66  int BLOCK_DIM_Y = 1,

67  int BLOCK_DIM_Z = 1,

68  int ARCH = HIPCUB_ARCH /* ignored */

69 >

70 class BlockHistogram

71  : private ::rocprim::block_histogram<

72  T,

73  BLOCK_DIM_X,

74  ITEMS_PER_THREAD,

75  BINS,

76  static_cast<::rocprim::block_histogram_algorithm>(ALGORITHM),

77  BLOCK_DIM_Y,

78  BLOCK_DIM_Z

79  >

80 {

81  static_assert(

82  BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z > 0,

83  "BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z must be greater than 0"

84  );

85 

86  using base_type =

87  typename ::rocprim::block_histogram<

88  T,

89  BLOCK_DIM_X,

90  ITEMS_PER_THREAD,

91  BINS,

92  static_cast<::rocprim::block_histogram_algorithm>(ALGORITHM),

93  BLOCK_DIM_Y,

94  BLOCK_DIM_Z

95  >;

96 

97  // Reference to temporary storage (usually shared memory)

98  typename base_type::storage_type& temp_storage_;

99 

100 public:

101  using TempStorage = typename base_type::storage_type;

102 

103  HIPCUB_DEVICE inline

104  BlockHistogram() : temp_storage_(private_storage())

105  {

106  }

107 

108  HIPCUB_DEVICE inline

109  BlockHistogram(TempStorage& temp_storage) : temp_storage_(temp_storage)

110  {

111  }

112 

113  template<class CounterT>

114  HIPCUB_DEVICE inline

115  void InitHistogram(CounterT histogram[BINS])

116  {

117  base_type::init_histogram(histogram);

118  }

119 

120  template<class CounterT>

121  HIPCUB_DEVICE inline

122  void Composite(T (&items)[ITEMS_PER_THREAD],

123  CounterT histogram[BINS])

124  {

125  base_type::composite(items, histogram, temp_storage_);

126  }

127 

128  template<class CounterT>

129  HIPCUB_DEVICE inline

130  void Histogram(T (&items)[ITEMS_PER_THREAD],

131  CounterT histogram[BINS])

132  {

133  base_type::init_histogram(histogram);

134  CTA_SYNC();

135  base_type::composite(items, histogram, temp_storage_);

136  }

137 

138 private:

139  HIPCUB_DEVICE inline

140  TempStorage& private_storage()

141  {

142  HIPCUB_SHARED_MEMORY TempStorage private_storage;

143  return private_storage;

144  }

145 };

146 

147 END_HIPCUB_NAMESPACE

148 

149 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_HISTOGRAM_HPP_

hipcub::BlockHistogram

Definition: block_histogram.hpp:80