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

• home
• docs
• checkouts
• readthedocs.org
• user_builds
• advanced-micro-devices-hipcub
• checkouts
• docs-5.7.0
• hipcub
• include
• hipcub
• backend
• cub
• 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-2023, 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_CUB_DEVICE_DEVICE_SCAN_HPP_

31 #define HIPCUB_CUB_DEVICE_DEVICE_SCAN_HPP_

32 

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

34 

35 #include <cub/device/device_scan.cuh>

36 

37 BEGIN_HIPCUB_NAMESPACE

38 

39 class DeviceScan

40 {

41 public:

42  template <

43  typename InputIteratorT,

44  typename OutputIteratorT

45  >

46  HIPCUB_RUNTIME_FUNCTION static

47  hipError_t InclusiveSum(void *d_temp_storage,

48  size_t &temp_storage_bytes,

49  InputIteratorT d_in,

50  OutputIteratorT d_out,

51  int num_items,

52  hipStream_t stream = 0,

53  bool debug_synchronous = false)

54  {

55  (void)debug_synchronous;

56  return hipCUDAErrorTohipError(::cub::DeviceScan::InclusiveSum(d_temp_storage,

57  temp_storage_bytes,

58  d_in,

59  d_out,

60  num_items,

61  stream));

62  }

63 

64  template <

65  typename InputIteratorT,

66  typename OutputIteratorT,

67  typename ScanOpT

68  >

69  HIPCUB_RUNTIME_FUNCTION static

70  hipError_t InclusiveScan(void *d_temp_storage,

71  size_t &temp_storage_bytes,

72  InputIteratorT d_in,

73  OutputIteratorT d_out,

74  ScanOpT scan_op,

75  int num_items,

76  hipStream_t stream = 0,

77  bool debug_synchronous = false)

78  {

79  (void)debug_synchronous;

80  return hipCUDAErrorTohipError(::cub::DeviceScan::InclusiveScan(d_temp_storage,

81  temp_storage_bytes,

82  d_in,

83  d_out,

84  scan_op,

85  num_items,

86  stream));

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  int num_items,

99  hipStream_t stream = 0,

100  bool debug_synchronous = false)

101  {

102  (void)debug_synchronous;

103  return hipCUDAErrorTohipError(::cub::DeviceScan::ExclusiveSum(d_temp_storage,

104  temp_storage_bytes,

105  d_in,

106  d_out,

107  num_items,

108  stream));

109  }

110 

111  template <

112  typename InputIteratorT,

113  typename OutputIteratorT,

114  typename ScanOpT,

115  typename InitValueT

116  >

117  HIPCUB_RUNTIME_FUNCTION static

118  hipError_t ExclusiveScan(void *d_temp_storage,

119  size_t &temp_storage_bytes,

120  InputIteratorT d_in,

121  OutputIteratorT d_out,

122  ScanOpT scan_op,

123  InitValueT init_value,

124  int num_items,

125  hipStream_t stream = 0,

126  bool debug_synchronous = false)

127  {

128  (void)debug_synchronous;

129  return hipCUDAErrorTohipError(::cub::DeviceScan::ExclusiveScan(d_temp_storage,

130  temp_storage_bytes,

131  d_in,

132  d_out,

133  scan_op,

134  init_value,

135  num_items,

136  stream));

137  }

138 

139  template <

140  typename InputIteratorT,

141  typename OutputIteratorT,

142  typename ScanOpT,

143  typename InitValueT,

144  typename InitValueIterT = InitValueT*

145  >

146  HIPCUB_RUNTIME_FUNCTION static

147  hipError_t ExclusiveScan(void *d_temp_storage,

148  size_t &temp_storage_bytes,

149  InputIteratorT d_in,

150  OutputIteratorT d_out,

151  ScanOpT scan_op,

152  FutureValue<InitValueT, InitValueIterT> init_value,

153  int num_items,

154  hipStream_t stream = 0,

155  bool debug_synchronous = false)

156  {

157  (void)debug_synchronous;

158  return hipCUDAErrorTohipError(::cub::DeviceScan::ExclusiveScan(d_temp_storage,

159  temp_storage_bytes,

160  d_in,

161  d_out,

162  scan_op,

163  init_value,

164  num_items,

165  stream));

166  }

167 

168  template <

169  typename KeysInputIteratorT,

170  typename ValuesInputIteratorT,

171  typename ValuesOutputIteratorT,

172  typename EqualityOpT = ::hipcub::Equality

173  >

174  HIPCUB_RUNTIME_FUNCTION static

175  hipError_t ExclusiveSumByKey(void *d_temp_storage,

176  size_t &temp_storage_bytes,

177  KeysInputIteratorT d_keys_in,

178  ValuesInputIteratorT d_values_in,

179  ValuesOutputIteratorT d_values_out,

180  int num_items,

181  EqualityOpT equality_op = EqualityOpT(),

182  hipStream_t stream = 0,

183  bool debug_synchronous = false)

184  {

185  (void)debug_synchronous;

186  return hipCUDAErrorTohipError(::cub::DeviceScan::ExclusiveSumByKey(d_temp_storage,

187  temp_storage_bytes,

188  d_keys_in,

189  d_values_in,

190  d_values_out,

191  num_items,

192  equality_op,

193  stream));

194  }

195 

196  template <

197  typename KeysInputIteratorT,

198  typename ValuesInputIteratorT,

199  typename ValuesOutputIteratorT,

200  typename ScanOpT,

201  typename InitValueT,

202  typename EqualityOpT = ::hipcub::Equality

203  >

204  HIPCUB_RUNTIME_FUNCTION static

205  hipError_t ExclusiveScanByKey(void *d_temp_storage,

206  size_t &temp_storage_bytes,

207  KeysInputIteratorT d_keys_in,

208  ValuesInputIteratorT d_values_in,

209  ValuesOutputIteratorT d_values_out,

210  ScanOpT scan_op,

211  InitValueT init_value,

212  int num_items,

213  EqualityOpT equality_op = EqualityOpT(),

214  hipStream_t stream = 0,

215  bool debug_synchronous = false)

216  {

217  (void)debug_synchronous;

218  return hipCUDAErrorTohipError(::cub::DeviceScan::ExclusiveScanByKey(d_temp_storage,

219  temp_storage_bytes,

220  d_keys_in,

221  d_values_in,

222  d_values_out,

223  scan_op,

224  init_value,

225  num_items,

226  equality_op,

227  stream));

228  }

229 

230  template <

231  typename KeysInputIteratorT,

232  typename ValuesInputIteratorT,

233  typename ValuesOutputIteratorT,

234  typename EqualityOpT = ::hipcub::Equality

235  >

236  HIPCUB_RUNTIME_FUNCTION static

237  hipError_t InclusiveSumByKey(void *d_temp_storage,

238  size_t &temp_storage_bytes,

239  KeysInputIteratorT d_keys_in,

240  ValuesInputIteratorT d_values_in,

241  ValuesOutputIteratorT d_values_out,

242  int num_items,

243  EqualityOpT equality_op = EqualityOpT(),

244  hipStream_t stream = 0,

245  bool debug_synchronous = false)

246  {

247  (void)debug_synchronous;

248  return hipCUDAErrorTohipError(::cub::DeviceScan::InclusiveSumByKey(d_temp_storage,

249  temp_storage_bytes,

250  d_keys_in,

251  d_values_in,

252  d_values_out,

253  num_items,

254  equality_op,

255  stream));

256  }

257 

258  template <

259  typename KeysInputIteratorT,

260  typename ValuesInputIteratorT,

261  typename ValuesOutputIteratorT,

262  typename ScanOpT,

263  typename EqualityOpT = ::hipcub::Equality

264  >

265  HIPCUB_RUNTIME_FUNCTION static

266  hipError_t InclusiveScanByKey(void *d_temp_storage,

267  size_t &temp_storage_bytes,

268  KeysInputIteratorT d_keys_in,

269  ValuesInputIteratorT d_values_in,

270  ValuesOutputIteratorT d_values_out,

271  ScanOpT scan_op,

272  int num_items,

273  EqualityOpT equality_op = EqualityOpT(),

274  hipStream_t stream = 0,

275  bool debug_synchronous = false)

276  {

277  (void)debug_synchronous;

278  return hipCUDAErrorTohipError(::cub::DeviceScan::InclusiveScanByKey(d_temp_storage,

279  temp_storage_bytes,

280  d_keys_in,

281  d_values_in,

282  d_values_out,

283  scan_op,

284  num_items,

285  equality_op,

286  stream));

287  }

288 };

289 

290 END_HIPCUB_NAMESPACE

291 

292 #endif // HIPCUB_CUB_DEVICE_DEVICE_SCAN_HPP_

hipcub::DeviceScan

Definition: device_scan.hpp:40

hipcub::Equality

Definition: thread_operators.hpp:40