/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.0/hipcub/include/hipcub/backend/rocprim/device/device_scan.hpp Source File

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

device_scan.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_DEVICE_DEVICE_SCAN_HPP_

31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_

32 

33 #include <iostream>

34 #include "../../../config.hpp"

35 

36 #include "../thread/thread_operators.hpp"

37 

38 #include <rocprim/device/device_scan.hpp>

39 #include <rocprim/device/device_scan_by_key.hpp>

40 

41 BEGIN_HIPCUB_NAMESPACE

42 

43 class DeviceScan

44 {

45 public:

46  template <

47  typename InputIteratorT,

48  typename OutputIteratorT

49  >

50  HIPCUB_RUNTIME_FUNCTION static

51  hipError_t InclusiveSum(void *d_temp_storage,

52  size_t &temp_storage_bytes,

53  InputIteratorT d_in,

54  OutputIteratorT d_out,

55  size_t num_items,

56  hipStream_t stream = 0,

57  bool debug_synchronous = false)

58  {

59  return InclusiveScan(

60  d_temp_storage, temp_storage_bytes,

61  d_in, d_out, ::hipcub::Sum(), num_items,

62  stream, debug_synchronous

63  );

64  }

65 

66  template <

67  typename InputIteratorT,

68  typename OutputIteratorT,

69  typename ScanOpT

70  >

71  HIPCUB_RUNTIME_FUNCTION static

72  hipError_t InclusiveScan(void *d_temp_storage,

73  size_t &temp_storage_bytes,

74  InputIteratorT d_in,

75  OutputIteratorT d_out,

76  ScanOpT scan_op,

77  size_t num_items,

78  hipStream_t stream = 0,

79  bool debug_synchronous = false)

80  {

81  return ::rocprim::inclusive_scan(

82  d_temp_storage, temp_storage_bytes,

83  d_in, d_out, num_items,

84  scan_op,

85  stream, debug_synchronous

86  );

87  }

88 

89  template <

90  typename InputIteratorT,

91  typename OutputIteratorT

92  >

93  HIPCUB_RUNTIME_FUNCTION static

94  hipError_t ExclusiveSum(void *d_temp_storage,

95  size_t &temp_storage_bytes,

96  InputIteratorT d_in,

97  OutputIteratorT d_out,

98  size_t num_items,

99  hipStream_t stream = 0,

100  bool debug_synchronous = false)

101  {

102  using T = typename std::iterator_traits<InputIteratorT>::value_type;

103  return ExclusiveScan(

104  d_temp_storage, temp_storage_bytes,

105  d_in, d_out, ::hipcub::Sum(), T(0), num_items,

106  stream, debug_synchronous

107  );

108  }

109 

110  template <

111  typename InputIteratorT,

112  typename OutputIteratorT,

113  typename ScanOpT,

114  typename InitValueT

115  >

116  HIPCUB_RUNTIME_FUNCTION static

117  hipError_t ExclusiveScan(void *d_temp_storage,

118  size_t &temp_storage_bytes,

119  InputIteratorT d_in,

120  OutputIteratorT d_out,

121  ScanOpT scan_op,

122  InitValueT init_value,

123  size_t num_items,

124  hipStream_t stream = 0,

125  bool debug_synchronous = false)

126  {

127  return ::rocprim::exclusive_scan(

128  d_temp_storage, temp_storage_bytes,

129  d_in, d_out, init_value, num_items,

130  scan_op,

131  stream, debug_synchronous

132  );

133  }

134 

135  template <

136  typename InputIteratorT,

137  typename OutputIteratorT,

138  typename ScanOpT,

139  typename InitValueT,

140  typename InitValueIterT = InitValueT*

141  >

142  HIPCUB_RUNTIME_FUNCTION static

143  hipError_t ExclusiveScan(void *d_temp_storage,

144  size_t &temp_storage_bytes,

145  InputIteratorT d_in,

146  OutputIteratorT d_out,

147  ScanOpT scan_op,

148  FutureValue<InitValueT, InitValueIterT> init_value,

149  int num_items,

150  hipStream_t stream = 0,

151  bool debug_synchronous = false)

152  {

153  return ::rocprim::exclusive_scan(

154  d_temp_storage, temp_storage_bytes,

155  d_in, d_out, init_value, num_items,

156  scan_op,

157  stream, debug_synchronous

158  );

159  }

160 

161  template <

162  typename KeysInputIteratorT,

163  typename ValuesInputIteratorT,

164  typename ValuesOutputIteratorT,

165  typename EqualityOpT = ::hipcub::Equality

166  >

167  HIPCUB_RUNTIME_FUNCTION static

168  hipError_t ExclusiveSumByKey(void *d_temp_storage,

169  size_t &temp_storage_bytes,

170  KeysInputIteratorT d_keys_in,

171  ValuesInputIteratorT d_values_in,

172  ValuesOutputIteratorT d_values_out,

173  int num_items,

174  EqualityOpT equality_op = EqualityOpT(),

175  hipStream_t stream = 0,

176  bool debug_synchronous = false)

177  {

178  using in_value_type = typename std::iterator_traits<ValuesInputIteratorT>::value_type;

179 

180  return ::rocprim::exclusive_scan_by_key(

181  d_temp_storage, temp_storage_bytes,

182  d_keys_in, d_values_in, d_values_out,

183  static_cast<in_value_type>(0), static_cast<size_t>(num_items),

184  ::hipcub::Sum(), equality_op, stream, debug_synchronous

185  );

186  }

187 

188  template <

189  typename KeysInputIteratorT,

190  typename ValuesInputIteratorT,

191  typename ValuesOutputIteratorT,

192  typename ScanOpT,

193  typename InitValueT,

194  typename EqualityOpT = ::hipcub::Equality

195  >

196  HIPCUB_RUNTIME_FUNCTION static

197  hipError_t ExclusiveScanByKey(void *d_temp_storage,

198  size_t &temp_storage_bytes,

199  KeysInputIteratorT d_keys_in,

200  ValuesInputIteratorT d_values_in,

201  ValuesOutputIteratorT d_values_out,

202  ScanOpT scan_op,

203  InitValueT init_value,

204  int num_items,

205  EqualityOpT equality_op = EqualityOpT(),

206  hipStream_t stream = 0,

207  bool debug_synchronous = false)

208  {

209  return ::rocprim::exclusive_scan_by_key(

210  d_temp_storage, temp_storage_bytes,

211  d_keys_in, d_values_in, d_values_out,

212  init_value, static_cast<size_t>(num_items),

213  scan_op, equality_op, stream, debug_synchronous

214  );

215  }

216 

217  template <

218  typename KeysInputIteratorT,

219  typename ValuesInputIteratorT,

220  typename ValuesOutputIteratorT,

221  typename EqualityOpT = ::hipcub::Equality

222  >

223  HIPCUB_RUNTIME_FUNCTION static

224  hipError_t InclusiveSumByKey(void *d_temp_storage,

225  size_t &temp_storage_bytes,

226  KeysInputIteratorT d_keys_in,

227  ValuesInputIteratorT d_values_in,

228  ValuesOutputIteratorT d_values_out,

229  int num_items,

230  EqualityOpT equality_op = EqualityOpT(),

231  hipStream_t stream = 0,

232  bool debug_synchronous = false)

233  {

234  return ::rocprim::inclusive_scan_by_key(

235  d_temp_storage, temp_storage_bytes,

236  d_keys_in, d_values_in, d_values_out,

237  static_cast<size_t>(num_items), ::hipcub::Sum(),

238  equality_op, stream, debug_synchronous

239  );

240  }

241 

242  template <

243  typename KeysInputIteratorT,

244  typename ValuesInputIteratorT,

245  typename ValuesOutputIteratorT,

246  typename ScanOpT,

247  typename EqualityOpT = ::hipcub::Equality

248  >

249  HIPCUB_RUNTIME_FUNCTION static

250  hipError_t InclusiveScanByKey(void *d_temp_storage,

251  size_t &temp_storage_bytes,

252  KeysInputIteratorT d_keys_in,

253  ValuesInputIteratorT d_values_in,

254  ValuesOutputIteratorT d_values_out,

255  ScanOpT scan_op,

256  int num_items,

257  EqualityOpT equality_op = EqualityOpT(),

258  hipStream_t stream = 0,

259  bool debug_synchronous = false)

260  {

261  return ::rocprim::inclusive_scan_by_key(

262  d_temp_storage, temp_storage_bytes,

263  d_keys_in, d_values_in, d_values_out,

264  static_cast<size_t>(num_items), scan_op,

265  equality_op, stream, debug_synchronous

266  );

267  }

268 };

269 

270 END_HIPCUB_NAMESPACE

271 

272 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_

hipcub::Equality

Definition: thread_operators.hpp:40

hipcub::Sum

Definition: thread_operators.hpp:76