/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/cub/device/device_segmented_radix_sort.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_segmented_radix_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-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_SEGMENTED_RADIX_SORT_HPP_

31 #define HIPCUB_CUB_DEVICE_DEVICE_SEGMENTED_RADIX_SORT_HPP_

32 

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

34 

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

36 

37 BEGIN_HIPCUB_NAMESPACE

38 

39 struct DeviceSegmentedRadixSort

40 {

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

42  HIPCUB_RUNTIME_FUNCTION static

43  hipError_t SortPairs(void * d_temp_storage,

44  size_t& temp_storage_bytes,

45  const KeyT * d_keys_in,

46  KeyT * d_keys_out,

47  const ValueT * d_values_in,

48  ValueT * d_values_out,

49  int num_items,

50  int num_segments,

51  OffsetIteratorT d_begin_offsets,

52  OffsetIteratorT d_end_offsets,

53  int begin_bit = 0,

54  int end_bit = sizeof(KeyT) * 8,

55  hipStream_t stream = 0,

56  bool debug_synchronous = false)

57  {

58  (void)debug_synchronous;

59  return hipCUDAErrorTohipError(::cub::DeviceSegmentedRadixSort::SortPairs(d_temp_storage,

60  temp_storage_bytes,

61  d_keys_in,

62  d_keys_out,

63  d_values_in,

64  d_values_out,

65  num_items,

66  num_segments,

67  d_begin_offsets,

68  d_end_offsets,

69  begin_bit,

70  end_bit,

71  stream));

72  }

73 

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

75  HIPCUB_RUNTIME_FUNCTION static

76  hipError_t SortPairs(void * d_temp_storage,

77  size_t& temp_storage_bytes,

78  DoubleBuffer<KeyT>& d_keys,

79  DoubleBuffer<ValueT>& d_values,

80  int num_items,

81  int num_segments,

82  OffsetIteratorT d_begin_offsets,

83  OffsetIteratorT d_end_offsets,

84  int begin_bit = 0,

85  int end_bit = sizeof(KeyT) * 8,

86  hipStream_t stream = 0,

87  bool debug_synchronous = false)

88  {

89  (void)debug_synchronous;

90  return hipCUDAErrorTohipError(::cub::DeviceSegmentedRadixSort::SortPairs(d_temp_storage,

91  temp_storage_bytes,

92  d_keys,

93  d_values,

94  num_items,

95  num_segments,

96  d_begin_offsets,

97  d_end_offsets,

98  begin_bit,

99  end_bit,

100  stream));

101  }

102 

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

104  HIPCUB_RUNTIME_FUNCTION static

105  hipError_t SortPairsDescending(void * d_temp_storage,

106  size_t& temp_storage_bytes,

107  const KeyT * d_keys_in,

108  KeyT * d_keys_out,

109  const ValueT * d_values_in,

110  ValueT * d_values_out,

111  int num_items,

112  int num_segments,

113  OffsetIteratorT d_begin_offsets,

114  OffsetIteratorT d_end_offsets,

115  int begin_bit = 0,

116  int end_bit = sizeof(KeyT) * 8,

117  hipStream_t stream = 0,

118  bool debug_synchronous = false)

119  {

120  (void)debug_synchronous;

121  return hipCUDAErrorTohipError(

122  ::cub::DeviceSegmentedRadixSort::SortPairsDescending(d_temp_storage,

123  temp_storage_bytes,

124  d_keys_in,

125  d_keys_out,

126  d_values_in,

127  d_values_out,

128  num_items,

129  num_segments,

130  d_begin_offsets,

131  d_end_offsets,

132  begin_bit,

133  end_bit,

134  stream));

135  }

136 

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

138  HIPCUB_RUNTIME_FUNCTION static

139  hipError_t SortPairsDescending(void * d_temp_storage,

140  size_t& temp_storage_bytes,

141  DoubleBuffer<KeyT>& d_keys,

142  DoubleBuffer<ValueT>& d_values,

143  int num_items,

144  int num_segments,

145  OffsetIteratorT d_begin_offsets,

146  OffsetIteratorT d_end_offsets,

147  int begin_bit = 0,

148  int end_bit = sizeof(KeyT) * 8,

149  hipStream_t stream = 0,

150  bool debug_synchronous = false)

151  {

152  (void)debug_synchronous;

153  return hipCUDAErrorTohipError(

154  ::cub::DeviceSegmentedRadixSort::SortPairsDescending(d_temp_storage,

155  temp_storage_bytes,

156  d_keys,

157  d_values,

158  num_items,

159  num_segments,

160  d_begin_offsets,

161  d_end_offsets,

162  begin_bit,

163  end_bit,

164  stream));

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  const KeyT * d_keys_in,

172  KeyT * d_keys_out,

173  int num_items,

174  int num_segments,

175  OffsetIteratorT d_begin_offsets,

176  OffsetIteratorT d_end_offsets,

177  int begin_bit = 0,

178  int end_bit = sizeof(KeyT) * 8,

179  hipStream_t stream = 0,

180  bool debug_synchronous = false)

181  {

182  (void)debug_synchronous;

183  return hipCUDAErrorTohipError(::cub::DeviceSegmentedRadixSort::SortKeys(d_temp_storage,

184  temp_storage_bytes,

185  d_keys_in,

186  d_keys_out,

187  num_items,

188  num_segments,

189  d_begin_offsets,

190  d_end_offsets,

191  begin_bit,

192  end_bit,

193  stream));

194  }

195 

196  template<typename KeyT, typename OffsetIteratorT>

197  HIPCUB_RUNTIME_FUNCTION static

198  hipError_t SortKeys(void * d_temp_storage,

199  size_t& temp_storage_bytes,

200  DoubleBuffer<KeyT>& d_keys,

201  int num_items,

202  int num_segments,

203  OffsetIteratorT d_begin_offsets,

204  OffsetIteratorT d_end_offsets,

205  int begin_bit = 0,

206  int end_bit = sizeof(KeyT) * 8,

207  hipStream_t stream = 0,

208  bool debug_synchronous = false)

209  {

210  (void)debug_synchronous;

211  return hipCUDAErrorTohipError(::cub::DeviceSegmentedRadixSort::SortKeys(d_temp_storage,

212  temp_storage_bytes,

213  d_keys,

214  num_items,

215  num_segments,

216  d_begin_offsets,

217  d_end_offsets,

218  begin_bit,

219  end_bit,

220  stream));

221  }

222 

223  template<typename KeyT, typename OffsetIteratorT>

224  HIPCUB_RUNTIME_FUNCTION static

225  hipError_t SortKeysDescending(void * d_temp_storage,

226  size_t& temp_storage_bytes,

227  const KeyT * d_keys_in,

228  KeyT * d_keys_out,

229  int num_items,

230  int num_segments,

231  OffsetIteratorT d_begin_offsets,

232  OffsetIteratorT d_end_offsets,

233  int begin_bit = 0,

234  int end_bit = sizeof(KeyT) * 8,

235  hipStream_t stream = 0,

236  bool debug_synchronous = false)

237  {

238  (void)debug_synchronous;

239  return hipCUDAErrorTohipError(

240  ::cub::DeviceSegmentedRadixSort::SortKeysDescending(d_temp_storage,

241  temp_storage_bytes,

242  d_keys_in,

243  d_keys_out,

244  num_items,

245  num_segments,

246  d_begin_offsets,

247  d_end_offsets,

248  begin_bit,

249  end_bit,

250  stream));

251  }

252 

253  template<typename KeyT, typename OffsetIteratorT>

254  HIPCUB_RUNTIME_FUNCTION static

255  hipError_t SortKeysDescending(void * d_temp_storage,

256  size_t& temp_storage_bytes,

257  DoubleBuffer<KeyT>& d_keys,

258  int num_items,

259  int num_segments,

260  OffsetIteratorT d_begin_offsets,

261  OffsetIteratorT d_end_offsets,

262  int begin_bit = 0,

263  int end_bit = sizeof(KeyT) * 8,

264  hipStream_t stream = 0,

265  bool debug_synchronous = false)

266  {

267  (void)debug_synchronous;

268  return hipCUDAErrorTohipError(

269  ::cub::DeviceSegmentedRadixSort::SortKeysDescending(d_temp_storage,

270  temp_storage_bytes,

271  d_keys,

272  num_items,

273  num_segments,

274  d_begin_offsets,

275  d_end_offsets,

276  begin_bit,

277  end_bit,

278  stream));

279  }

280 };

281 

282 END_HIPCUB_NAMESPACE

283 

284 #endif // HIPCUB_CUB_DEVICE_DEVICE_SEGMENTED_RADIX_SORT_HPP_

hipcub::DeviceSegmentedRadixSort

Definition: device_segmented_radix_sort.hpp:40

hipcub::DoubleBuffer

Definition: util_type.hpp:108