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

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

device_segmented_sort.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_SEGMENTED_SORT_HPP_

31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_SORT_HPP_

32 

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

34 

35 #include "../util_type.hpp"

36 

37 #include <rocprim/device/device_segmented_radix_sort.hpp>

38 

39 BEGIN_HIPCUB_NAMESPACE

40 

41 struct DeviceSegmentedSort

42 {

43  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

44  HIPCUB_RUNTIME_FUNCTION static

45  hipError_t SortPairs(void * d_temp_storage,

46  size_t& temp_storage_bytes,

47  const KeyT * d_keys_in,

48  KeyT * d_keys_out,

49  const ValueT * d_values_in,

50  ValueT * d_values_out,

51  int num_items,

52  int num_segments,

53  OffsetIteratorT d_begin_offsets,

54  OffsetIteratorT d_end_offsets,

55  hipStream_t stream = 0,

56  bool debug_synchronous = false)

57  {

58  return ::rocprim::segmented_radix_sort_pairs(

59  d_temp_storage, temp_storage_bytes,

60  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

61  num_segments, d_begin_offsets, d_end_offsets,

62  0, sizeof(KeyT) * 8,

63  stream, debug_synchronous

64  );

65  }

66 

67  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

68  HIPCUB_RUNTIME_FUNCTION static

69  hipError_t SortPairs(void * d_temp_storage,

70  size_t& temp_storage_bytes,

71  DoubleBuffer<KeyT>& d_keys,

72  DoubleBuffer<ValueT>& d_values,

73  int num_items,

74  int num_segments,

75  OffsetIteratorT d_begin_offsets,

76  OffsetIteratorT d_end_offsets,

77  hipStream_t stream = 0,

78  bool debug_synchronous = false)

79  {

80  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);

81  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);

82  hipError_t error = ::rocprim::segmented_radix_sort_pairs(

83  d_temp_storage, temp_storage_bytes,

84  d_keys_db, d_values_db, num_items,

85  num_segments, d_begin_offsets, d_end_offsets,

86  0, sizeof(KeyT) * 8,

87  stream, debug_synchronous

88  );

89  detail::update_double_buffer(d_keys, d_keys_db);

90  detail::update_double_buffer(d_values, d_values_db);

91  return error;

92  }

93 

94  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

95  HIPCUB_RUNTIME_FUNCTION static

96  hipError_t SortPairsDescending(void * d_temp_storage,

97  size_t& temp_storage_bytes,

98  const KeyT * d_keys_in,

99  KeyT * d_keys_out,

100  const ValueT * d_values_in,

101  ValueT * d_values_out,

102  int num_items,

103  int num_segments,

104  OffsetIteratorT d_begin_offsets,

105  OffsetIteratorT d_end_offsets,

106  hipStream_t stream = 0,

107  bool debug_synchronous = false)

108  {

109  return ::rocprim::segmented_radix_sort_pairs_desc(

110  d_temp_storage, temp_storage_bytes,

111  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

112  num_segments, d_begin_offsets, d_end_offsets,

113  0, sizeof(KeyT) * 8,

114  stream, debug_synchronous

115  );

116  }

117 

118  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

119  HIPCUB_RUNTIME_FUNCTION static

120  hipError_t SortPairsDescending(void * d_temp_storage,

121  size_t& temp_storage_bytes,

122  DoubleBuffer<KeyT>& d_keys,

123  DoubleBuffer<ValueT>& d_values,

124  int num_items,

125  int num_segments,

126  OffsetIteratorT d_begin_offsets,

127  OffsetIteratorT d_end_offsets,

128  hipStream_t stream = 0,

129  bool debug_synchronous = false)

130  {

131  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);

132  ::rocprim::double_buffer<ValueT> d_values_db = detail::to_double_buffer(d_values);

133  hipError_t error = ::rocprim::segmented_radix_sort_pairs_desc(

134  d_temp_storage, temp_storage_bytes,

135  d_keys_db, d_values_db, num_items,

136  num_segments, d_begin_offsets, d_end_offsets,

137  0, sizeof(KeyT) * 8,

138  stream, debug_synchronous

139  );

140  detail::update_double_buffer(d_keys, d_keys_db);

141  detail::update_double_buffer(d_values, d_values_db);

142  return error;

143  }

144 

145  template<typename KeyT, typename OffsetIteratorT>

146  HIPCUB_RUNTIME_FUNCTION static

147  hipError_t SortKeys(void * d_temp_storage,

148  size_t& temp_storage_bytes,

149  const KeyT * d_keys_in,

150  KeyT * d_keys_out,

151  int num_items,

152  int num_segments,

153  OffsetIteratorT d_begin_offsets,

154  OffsetIteratorT d_end_offsets,

155  hipStream_t stream = 0,

156  bool debug_synchronous = false)

157  {

158  return ::rocprim::segmented_radix_sort_keys(

159  d_temp_storage, temp_storage_bytes,

160  d_keys_in, d_keys_out, num_items,

161  num_segments, d_begin_offsets, d_end_offsets,

162  0, sizeof(KeyT) * 8,

163  stream, debug_synchronous

164  );

165  }

166 

167  template<typename KeyT, typename OffsetIteratorT>

168  HIPCUB_RUNTIME_FUNCTION static

169  hipError_t SortKeys(void * d_temp_storage,

170  size_t& temp_storage_bytes,

171  DoubleBuffer<KeyT>& d_keys,

172  int num_items,

173  int num_segments,

174  OffsetIteratorT d_begin_offsets,

175  OffsetIteratorT d_end_offsets,

176  hipStream_t stream = 0,

177  bool debug_synchronous = false)

178  {

179  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);

180  hipError_t error = ::rocprim::segmented_radix_sort_keys(

181  d_temp_storage, temp_storage_bytes,

182  d_keys_db, num_items,

183  num_segments, d_begin_offsets, d_end_offsets,

184  0, sizeof(KeyT) * 8,

185  stream, debug_synchronous

186  );

187  detail::update_double_buffer(d_keys, d_keys_db);

188  return error;

189  }

190 

191  template<typename KeyT, typename OffsetIteratorT>

192  HIPCUB_RUNTIME_FUNCTION static

193  hipError_t SortKeysDescending(void * d_temp_storage,

194  size_t& temp_storage_bytes,

195  const KeyT * d_keys_in,

196  KeyT * d_keys_out,

197  int num_items,

198  int num_segments,

199  OffsetIteratorT d_begin_offsets,

200  OffsetIteratorT d_end_offsets,

201  hipStream_t stream = 0,

202  bool debug_synchronous = false)

203  {

204  return ::rocprim::segmented_radix_sort_keys_desc(

205  d_temp_storage, temp_storage_bytes,

206  d_keys_in, d_keys_out, num_items,

207  num_segments, d_begin_offsets, d_end_offsets,

208  0, sizeof(KeyT) * 8,

209  stream, debug_synchronous

210  );

211  }

212 

213  template<typename KeyT, typename OffsetIteratorT>

214  HIPCUB_RUNTIME_FUNCTION static

215  hipError_t SortKeysDescending(void * d_temp_storage,

216  size_t& temp_storage_bytes,

217  DoubleBuffer<KeyT>& d_keys,

218  int num_items,

219  int num_segments,

220  OffsetIteratorT d_begin_offsets,

221  OffsetIteratorT d_end_offsets,

222  hipStream_t stream = 0,

223  bool debug_synchronous = false)

224  {

225  ::rocprim::double_buffer<KeyT> d_keys_db = detail::to_double_buffer(d_keys);

226  hipError_t error = ::rocprim::segmented_radix_sort_keys_desc(

227  d_temp_storage, temp_storage_bytes,

228  d_keys_db, num_items,

229  num_segments, d_begin_offsets, d_end_offsets,

230  0, sizeof(KeyT) * 8,

231  stream, debug_synchronous

232  );

233  detail::update_double_buffer(d_keys, d_keys_db);

234  return error;

235  }

236 

237  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

238  HIPCUB_RUNTIME_FUNCTION static

239  hipError_t StableSortPairs(void * d_temp_storage,

240  size_t& temp_storage_bytes,

241  const KeyT * d_keys_in,

242  KeyT * d_keys_out,

243  const ValueT * d_values_in,

244  ValueT * d_values_out,

245  int num_items,

246  int num_segments,

247  OffsetIteratorT d_begin_offsets,

248  OffsetIteratorT d_end_offsets,

249  hipStream_t stream = 0,

250  bool debug_synchronous = false)

251  {

252  return SortPairs(

253  d_temp_storage, temp_storage_bytes,

254  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

255  num_segments, d_begin_offsets, d_end_offsets,

256  stream, debug_synchronous

257  );

258  }

259 

260  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

261  HIPCUB_RUNTIME_FUNCTION static

262  hipError_t StableSortPairs(void * d_temp_storage,

263  size_t& temp_storage_bytes,

264  DoubleBuffer<KeyT>& d_keys,

265  DoubleBuffer<ValueT>& d_values,

266  int num_items,

267  int num_segments,

268  OffsetIteratorT d_begin_offsets,

269  OffsetIteratorT d_end_offsets,

270  hipStream_t stream = 0,

271  bool debug_synchronous = false)

272  {

273  return SortPairs(

274  d_temp_storage, temp_storage_bytes,

275  d_keys, d_values, num_items,

276  num_segments, d_begin_offsets, d_end_offsets,

277  stream, debug_synchronous

278  );

279  }

280 

281  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

282  HIPCUB_RUNTIME_FUNCTION static

283  hipError_t StableSortPairsDescending(void * d_temp_storage,

284  size_t& temp_storage_bytes,

285  const KeyT * d_keys_in,

286  KeyT * d_keys_out,

287  const ValueT * d_values_in,

288  ValueT * d_values_out,

289  int num_items,

290  int num_segments,

291  OffsetIteratorT d_begin_offsets,

292  OffsetIteratorT d_end_offsets,

293  hipStream_t stream = 0,

294  bool debug_synchronous = false)

295  {

296  return SortPairsDescending(

297  d_temp_storage, temp_storage_bytes,

298  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

299  num_segments, d_begin_offsets, d_end_offsets,

300  stream, debug_synchronous

301  );

302  }

303 

304  template<typename KeyT, typename ValueT, typename OffsetIteratorT>

305  HIPCUB_RUNTIME_FUNCTION static

306  hipError_t StableSortPairsDescending(void * d_temp_storage,

307  size_t& temp_storage_bytes,

308  DoubleBuffer<KeyT>& d_keys,

309  DoubleBuffer<ValueT>& d_values,

310  int num_items,

311  int num_segments,

312  OffsetIteratorT d_begin_offsets,

313  OffsetIteratorT d_end_offsets,

314  hipStream_t stream = 0,

315  bool debug_synchronous = false)

316  {

317  return SortPairsDescending(

318  d_temp_storage, temp_storage_bytes,

319  d_keys, d_values, num_items,

320  num_segments, d_begin_offsets, d_end_offsets,

321  stream, debug_synchronous

322  );

323  }

324 

325  template<typename KeyT, typename OffsetIteratorT>

326  HIPCUB_RUNTIME_FUNCTION static

327  hipError_t StableSortKeys(void * d_temp_storage,

328  size_t& temp_storage_bytes,

329  const KeyT * d_keys_in,

330  KeyT * d_keys_out,

331  int num_items,

332  int num_segments,

333  OffsetIteratorT d_begin_offsets,

334  OffsetIteratorT d_end_offsets,

335  hipStream_t stream = 0,

336  bool debug_synchronous = false)

337  {

338  return SortKeys(

339  d_temp_storage, temp_storage_bytes,

340  d_keys_in, d_keys_out, num_items,

341  num_segments, d_begin_offsets, d_end_offsets,

342  stream, debug_synchronous

343  );

344  }

345 

346  template<typename KeyT, typename OffsetIteratorT>

347  HIPCUB_RUNTIME_FUNCTION static

348  hipError_t StableSortKeys(void * d_temp_storage,

349  size_t& temp_storage_bytes,

350  DoubleBuffer<KeyT>& d_keys,

351  int num_items,

352  int num_segments,

353  OffsetIteratorT d_begin_offsets,

354  OffsetIteratorT d_end_offsets,

355  hipStream_t stream = 0,

356  bool debug_synchronous = false)

357  {

358  return SortKeys(

359  d_temp_storage, temp_storage_bytes,

360  d_keys, num_items,

361  num_segments, d_begin_offsets, d_end_offsets,

362  stream, debug_synchronous

363  );

364  }

365 

366  template<typename KeyT, typename OffsetIteratorT>

367  HIPCUB_RUNTIME_FUNCTION static

368  hipError_t StableSortKeysDescending(void * d_temp_storage,

369  size_t& temp_storage_bytes,

370  const KeyT * d_keys_in,

371  KeyT * d_keys_out,

372  int num_items,

373  int num_segments,

374  OffsetIteratorT d_begin_offsets,

375  OffsetIteratorT d_end_offsets,

376  hipStream_t stream = 0,

377  bool debug_synchronous = false)

378  {

379  return SortKeysDescending(

380  d_temp_storage, temp_storage_bytes,

381  d_keys_in, d_keys_out, num_items,

382  num_segments, d_begin_offsets, d_end_offsets,

383  stream, debug_synchronous

384  );

385  }

386 

387  template<typename KeyT, typename OffsetIteratorT>

388  HIPCUB_RUNTIME_FUNCTION static

389  hipError_t StableSortKeysDescending(void * d_temp_storage,

390  size_t& temp_storage_bytes,

391  DoubleBuffer<KeyT>& d_keys,

392  int num_items,

393  int num_segments,

394  OffsetIteratorT d_begin_offsets,

395  OffsetIteratorT d_end_offsets,

396  hipStream_t stream = 0,

397  bool debug_synchronous = false)

398  {

399  return SortKeysDescending(

400  d_temp_storage, temp_storage_bytes,

401  d_keys, num_items,

402  num_segments, d_begin_offsets, d_end_offsets,

403  stream, debug_synchronous

404  );

405  }

406 };

407 

408 END_HIPCUB_NAMESPACE

409 

410 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_SORT_HPP_