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

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

block_discontinuity.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-2020, 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_DISCONTINUITY_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_DISCONTINUITY_HPP_

32 

33 #include "../../../config.hpp"

34 

35 #include <rocprim/block/block_discontinuity.hpp>

36 

37 BEGIN_HIPCUB_NAMESPACE

38 

39 template<

40  typename T,

41  int BLOCK_DIM_X,

42  int BLOCK_DIM_Y = 1,

43  int BLOCK_DIM_Z = 1,

44  int ARCH = HIPCUB_ARCH /* ignored */

45 >

46 class BlockDiscontinuity

47  : private ::rocprim::block_discontinuity<

48  T,

49  BLOCK_DIM_X,

50  BLOCK_DIM_Y,

51  BLOCK_DIM_Z

52  >

53 {

54  static_assert(

55  BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z > 0,

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

57  );

58 

59  using base_type =

60  typename ::rocprim::block_discontinuity<

61  T,

62  BLOCK_DIM_X,

63  BLOCK_DIM_Y,

64  BLOCK_DIM_Z

65  >;

66 

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

68  typename base_type::storage_type& temp_storage_;

69 

70 public:

71  using TempStorage = typename base_type::storage_type;

72 

73  HIPCUB_DEVICE inline

74  BlockDiscontinuity() : temp_storage_(private_storage())

75  {

76  }

77 

78  HIPCUB_DEVICE inline

79  BlockDiscontinuity(TempStorage& temp_storage) : temp_storage_(temp_storage)

80  {

81  }

82 

83  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

84  HIPCUB_DEVICE inline

85  void FlagHeads(FlagT (&head_flags)[ITEMS_PER_THREAD],

86  T (&input)[ITEMS_PER_THREAD],

87  FlagOp flag_op)

88  {

89  base_type::flag_heads(head_flags, input, flag_op, temp_storage_);

90  }

91 

92  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

93  HIPCUB_DEVICE inline

94  void FlagHeads(FlagT (&head_flags)[ITEMS_PER_THREAD],

95  T (&input)[ITEMS_PER_THREAD],

96  FlagOp flag_op,

97  T tile_predecessor_item)

98  {

99  base_type::flag_heads(head_flags, tile_predecessor_item, input, flag_op, temp_storage_);

100  }

101 

102  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

103  HIPCUB_DEVICE inline

104  void FlagTails(FlagT (&tail_flags)[ITEMS_PER_THREAD],

105  T (&input)[ITEMS_PER_THREAD],

106  FlagOp flag_op)

107  {

108  base_type::flag_tails(tail_flags, input, flag_op, temp_storage_);

109  }

110 

111  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

112  HIPCUB_DEVICE inline

113  void FlagTails(FlagT (&tail_flags)[ITEMS_PER_THREAD],

114  T (&input)[ITEMS_PER_THREAD],

115  FlagOp flag_op,

116  T tile_successor_item)

117  {

118  base_type::flag_tails(tail_flags, tile_successor_item, input, flag_op, temp_storage_);

119  }

120 

121  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

122  HIPCUB_DEVICE inline

123  void FlagHeadsAndTails(FlagT (&head_flags)[ITEMS_PER_THREAD],

124  FlagT (&tail_flags)[ITEMS_PER_THREAD],

125  T (&input)[ITEMS_PER_THREAD],

126  FlagOp flag_op)

127  {

128  base_type::flag_heads_and_tails(

129  head_flags, tail_flags, input,

130  flag_op, temp_storage_

131  );

132  }

133 

134  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

135  HIPCUB_DEVICE inline

136  void FlagHeadsAndTails(FlagT (&head_flags)[ITEMS_PER_THREAD],

137  FlagT (&tail_flags)[ITEMS_PER_THREAD],

138  T tile_successor_item,

139  T (&input)[ITEMS_PER_THREAD],

140  FlagOp flag_op)

141  {

142  base_type::flag_heads_and_tails(

143  head_flags, tail_flags, tile_successor_item, input,

144  flag_op, temp_storage_

145  );

146  }

147 

148  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

149  HIPCUB_DEVICE inline

150  void FlagHeadsAndTails(FlagT (&head_flags)[ITEMS_PER_THREAD],

151  T tile_predecessor_item,

152  FlagT (&tail_flags)[ITEMS_PER_THREAD],

153  T (&input)[ITEMS_PER_THREAD],

154  FlagOp flag_op)

155  {

156  base_type::flag_heads_and_tails(

157  head_flags, tile_predecessor_item, tail_flags, input,

158  flag_op, temp_storage_

159  );

160  }

161 

162  template<int ITEMS_PER_THREAD, typename FlagT, typename FlagOp>

163  HIPCUB_DEVICE inline

164  void FlagHeadsAndTails(FlagT (&head_flags)[ITEMS_PER_THREAD],

165  T tile_predecessor_item,

166  FlagT (&tail_flags)[ITEMS_PER_THREAD],

167  T tile_successor_item,

168  T (&input)[ITEMS_PER_THREAD],

169  FlagOp flag_op)

170  {

171  base_type::flag_heads_and_tails(

172  head_flags, tile_predecessor_item, tail_flags, tile_successor_item, input,

173  flag_op, temp_storage_

174  );

175  }

176 

177 private:

178  HIPCUB_DEVICE inline

179  TempStorage& private_storage()

180  {

181  HIPCUB_SHARED_MEMORY TempStorage private_storage;

182  return private_storage;

183  }

184 };

185 

186 END_HIPCUB_NAMESPACE

187 

188 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_DISCONTINUITY_HPP_

hipcub::BlockDiscontinuity

Definition: block_discontinuity.hpp:53