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

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

block_adjacent_difference.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_ADJACENT_DIFFERENCE_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_ADJACENT_DIFFERENCE_HPP_

32 

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

34 

35 #include <rocprim/block/block_adjacent_difference.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 BlockAdjacentDifference

47  : private ::rocprim::block_adjacent_difference<

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_adjacent_difference<

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  BlockAdjacentDifference() : temp_storage_(private_storage())

75  {

76  }

77 

78  HIPCUB_DEVICE inline

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

80  {

81  }

82 

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

84  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

85  HIPCUB_DEVICE inline

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

87  T (&input)[ITEMS_PER_THREAD],

88  FlagOp flag_op)

89  {

90  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

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

92  HIPCUB_CLANG_SUPPRESS_WARNING_POP

93  }

94 

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

96  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

97  HIPCUB_DEVICE inline

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

99  T (&input)[ITEMS_PER_THREAD],

100  FlagOp flag_op,

101  T tile_predecessor_item)

102  {

103  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

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

105  HIPCUB_CLANG_SUPPRESS_WARNING_POP

106  }

107 

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

109  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

110  HIPCUB_DEVICE inline

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

112  T (&input)[ITEMS_PER_THREAD],

113  FlagOp flag_op)

114  {

115  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

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

117  HIPCUB_CLANG_SUPPRESS_WARNING_POP

118  }

119 

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

121  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

122  HIPCUB_DEVICE inline

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

124  T (&input)[ITEMS_PER_THREAD],

125  FlagOp flag_op,

126  T tile_successor_item)

127  {

128  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

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

130  HIPCUB_CLANG_SUPPRESS_WARNING_POP

131  }

132 

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

134  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

135  HIPCUB_DEVICE inline

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

137  FlagT (&tail_flags)[ITEMS_PER_THREAD],

138  T (&input)[ITEMS_PER_THREAD],

139  FlagOp flag_op)

140  {

141  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

142  base_type::flag_heads_and_tails(

143  head_flags, tail_flags, input,

144  flag_op, temp_storage_

145  );

146  HIPCUB_CLANG_SUPPRESS_WARNING_POP

147  }

148 

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

150  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

151  HIPCUB_DEVICE inline

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

153  FlagT (&tail_flags)[ITEMS_PER_THREAD],

154  T tile_successor_item,

155  T (&input)[ITEMS_PER_THREAD],

156  FlagOp flag_op)

157  {

158  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

159  base_type::flag_heads_and_tails(

160  head_flags, tail_flags, tile_successor_item, input,

161  flag_op, temp_storage_

162  );

163  HIPCUB_CLANG_SUPPRESS_WARNING_POP

164  }

165 

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

167  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

168  HIPCUB_DEVICE inline

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

170  T tile_predecessor_item,

171  FlagT (&tail_flags)[ITEMS_PER_THREAD],

172  T (&input)[ITEMS_PER_THREAD],

173  FlagOp flag_op)

174  {

175  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

176  base_type::flag_heads_and_tails(

177  head_flags, tile_predecessor_item, tail_flags, input,

178  flag_op, temp_storage_

179  );

180  HIPCUB_CLANG_SUPPRESS_WARNING_POP

181  }

182 

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

184  [[deprecated("The Flags API of BlockAdjacentDifference is deprecated.")]]

185  HIPCUB_DEVICE inline

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

187  T tile_predecessor_item,

188  FlagT (&tail_flags)[ITEMS_PER_THREAD],

189  T tile_successor_item,

190  T (&input)[ITEMS_PER_THREAD],

191  FlagOp flag_op)

192  {

193  HIPCUB_CLANG_SUPPRESS_WARNING_WITH_PUSH("-Wdeprecated")

194  base_type::flag_heads_and_tails(

195  head_flags, tile_predecessor_item, tail_flags, tile_successor_item, input,

196  flag_op, temp_storage_

197  );

198  HIPCUB_CLANG_SUPPRESS_WARNING_POP

199  }

200 

201  template <int ITEMS_PER_THREAD, typename OutputType, typename DifferenceOpT>

202  HIPCUB_DEVICE inline

203  void SubtractLeft(T (&input)[ITEMS_PER_THREAD],

204  OutputType (&output)[ITEMS_PER_THREAD],

205  DifferenceOpT difference_op)

206  {

207  base_type::subtract_left(

208  input, output, difference_op, temp_storage_

209  );

210  }

211 

212  template <int ITEMS_PER_THREAD, typename OutputT, typename DifferenceOpT>

213  HIPCUB_DEVICE inline

214  void SubtractLeft(T (&input)[ITEMS_PER_THREAD],

215  OutputT (&output)[ITEMS_PER_THREAD],

216  DifferenceOpT difference_op,

217  T tile_predecessor_item)

218  {

219  base_type::subtract_left(

220  input, output, difference_op, tile_predecessor_item, temp_storage_

221  );

222  }

223 

224  template <int ITEMS_PER_THREAD, typename OutputType, typename DifferenceOpT>

225  HIPCUB_DEVICE inline

226  void SubtractLeftPartialTile(T (&input)[ITEMS_PER_THREAD],

227  OutputType (&output)[ITEMS_PER_THREAD],

228  DifferenceOpT difference_op,

229  int valid_items)

230  {

231  base_type::subtract_left_partial(

232  input, output, difference_op, valid_items, temp_storage_

233  );

234  }

235 

236  template<int ITEMS_PER_THREAD, typename OutputType, typename DifferenceOpT>

237  HIPCUB_DEVICE inline void SubtractLeftPartialTile(T (&input)[ITEMS_PER_THREAD],

238  OutputType (&output)[ITEMS_PER_THREAD],

239  DifferenceOpT difference_op,

240  int valid_items,

241  T tile_predecessor_item)

242  {

243  base_type::subtract_left_partial(input,

244  output,

245  difference_op,

246  tile_predecessor_item,

247  valid_items,

248  temp_storage_);

249  }

250 

251  template <int ITEMS_PER_THREAD, typename OutputT, typename DifferenceOpT>

252  HIPCUB_DEVICE inline

253  void SubtractRight(T (&input)[ITEMS_PER_THREAD],

254  OutputT (&output)[ITEMS_PER_THREAD],

255  DifferenceOpT difference_op)

256  {

257  base_type::subtract_right(

258  input, output, difference_op, temp_storage_

259  );

260  }

261 

262  template <int ITEMS_PER_THREAD, typename OutputT, typename DifferenceOpT>

263  HIPCUB_DEVICE inline

264  void SubtractRight(T (&input)[ITEMS_PER_THREAD],

265  OutputT (&output)[ITEMS_PER_THREAD],

266  DifferenceOpT difference_op,

267  T tile_successor_item)

268  {

269  base_type::subtract_right(

270  input, output, difference_op, tile_successor_item, temp_storage_

271  );

272  }

273 

274  template <int ITEMS_PER_THREAD, typename OutputT, typename DifferenceOpT>

275  HIPCUB_DEVICE inline

276  void SubtractRightPartialTile(T (&input)[ITEMS_PER_THREAD],

277  OutputT (&output)[ITEMS_PER_THREAD],

278  DifferenceOpT difference_op,

279  int valid_items)

280  {

281  base_type::subtract_right_partial(

282  input, output, difference_op, valid_items, temp_storage_

283  );

284  }

285 

286 private:

287  HIPCUB_DEVICE inline

288  TempStorage& private_storage()

289  {

290  HIPCUB_SHARED_MEMORY TempStorage private_storage;

291  return private_storage;

292  }

293 };

294 

295 END_HIPCUB_NAMESPACE

296 

297 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_ADJACENT_DIFFERENCE_HPP_

hipcub::BlockAdjacentDifference

Definition: block_adjacent_difference.hpp:53