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

• home
• docs
• checkouts
• readthedocs.org
• user_builds
• advanced-micro-devices-hipcub
• checkouts
• docs-5.5.1
• hipcub
• include
• hipcub
• backend
• rocprim
• 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-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_RADIX_SORT_HPP_

31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_RADIX_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 DeviceSegmentedRadixSort

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  int begin_bit = 0,

56  int end_bit = sizeof(KeyT) * 8,

57  hipStream_t stream = 0,

58  bool debug_synchronous = false)

59  {

60  return ::rocprim::segmented_radix_sort_pairs(

61  d_temp_storage, temp_storage_bytes,

62  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

63  num_segments, d_begin_offsets, d_end_offsets,

64  begin_bit, end_bit,

65  stream, debug_synchronous

66  );

67  }

68 

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

70  HIPCUB_RUNTIME_FUNCTION static

71  hipError_t SortPairs(void * d_temp_storage,

72  size_t& temp_storage_bytes,

73  DoubleBuffer<KeyT>& d_keys,

74  DoubleBuffer<ValueT>& d_values,

75  int num_items,

76  int num_segments,

77  OffsetIteratorT d_begin_offsets,

78  OffsetIteratorT d_end_offsets,

79  int begin_bit = 0,

80  int end_bit = sizeof(KeyT) * 8,

81  hipStream_t stream = 0,

82  bool debug_synchronous = false)

83  {

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

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

86  hipError_t error = ::rocprim::segmented_radix_sort_pairs(

87  d_temp_storage, temp_storage_bytes,

88  d_keys_db, d_values_db, num_items,

89  num_segments, d_begin_offsets, d_end_offsets,

90  begin_bit, end_bit,

91  stream, debug_synchronous

92  );

93  detail::update_double_buffer(d_keys, d_keys_db);

94  detail::update_double_buffer(d_values, d_values_db);

95  return error;

96  }

97 

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

99  HIPCUB_RUNTIME_FUNCTION static

100  hipError_t SortPairsDescending(void * d_temp_storage,

101  size_t& temp_storage_bytes,

102  const KeyT * d_keys_in,

103  KeyT * d_keys_out,

104  const ValueT * d_values_in,

105  ValueT * d_values_out,

106  int num_items,

107  int num_segments,

108  OffsetIteratorT d_begin_offsets,

109  OffsetIteratorT d_end_offsets,

110  int begin_bit = 0,

111  int end_bit = sizeof(KeyT) * 8,

112  hipStream_t stream = 0,

113  bool debug_synchronous = false)

114  {

115  return ::rocprim::segmented_radix_sort_pairs_desc(

116  d_temp_storage, temp_storage_bytes,

117  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

118  num_segments, d_begin_offsets, d_end_offsets,

119  begin_bit, end_bit,

120  stream, debug_synchronous

121  );

122  }

123 

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

125  HIPCUB_RUNTIME_FUNCTION static

126  hipError_t SortPairsDescending(void * d_temp_storage,

127  size_t& temp_storage_bytes,

128  DoubleBuffer<KeyT>& d_keys,

129  DoubleBuffer<ValueT>& d_values,

130  int num_items,

131  int num_segments,

132  OffsetIteratorT d_begin_offsets,

133  OffsetIteratorT d_end_offsets,

134  int begin_bit = 0,

135  int end_bit = sizeof(KeyT) * 8,

136  hipStream_t stream = 0,

137  bool debug_synchronous = false)

138  {

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

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

141  hipError_t error = ::rocprim::segmented_radix_sort_pairs_desc(

142  d_temp_storage, temp_storage_bytes,

143  d_keys_db, d_values_db, num_items,

144  num_segments, d_begin_offsets, d_end_offsets,

145  begin_bit, end_bit,

146  stream, debug_synchronous

147  );

148  detail::update_double_buffer(d_keys, d_keys_db);

149  detail::update_double_buffer(d_values, d_values_db);

150  return error;

151  }

152 

153  template<typename KeyT, typename OffsetIteratorT>

154  HIPCUB_RUNTIME_FUNCTION static

155  hipError_t SortKeys(void * d_temp_storage,

156  size_t& temp_storage_bytes,

157  const KeyT * d_keys_in,

158  KeyT * d_keys_out,

159  int num_items,

160  int num_segments,

161  OffsetIteratorT d_begin_offsets,

162  OffsetIteratorT d_end_offsets,

163  int begin_bit = 0,

164  int end_bit = sizeof(KeyT) * 8,

165  hipStream_t stream = 0,

166  bool debug_synchronous = false)

167  {

168  return ::rocprim::segmented_radix_sort_keys(

169  d_temp_storage, temp_storage_bytes,

170  d_keys_in, d_keys_out, num_items,

171  num_segments, d_begin_offsets, d_end_offsets,

172  begin_bit, end_bit,

173  stream, debug_synchronous

174  );

175  }

176 

177  template<typename KeyT, typename OffsetIteratorT>

178  HIPCUB_RUNTIME_FUNCTION static

179  hipError_t SortKeys(void * d_temp_storage,

180  size_t& temp_storage_bytes,

181  DoubleBuffer<KeyT>& d_keys,

182  int num_items,

183  int num_segments,

184  OffsetIteratorT d_begin_offsets,

185  OffsetIteratorT d_end_offsets,

186  int begin_bit = 0,

187  int end_bit = sizeof(KeyT) * 8,

188  hipStream_t stream = 0,

189  bool debug_synchronous = false)

190  {

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

192  hipError_t error = ::rocprim::segmented_radix_sort_keys(

193  d_temp_storage, temp_storage_bytes,

194  d_keys_db, num_items,

195  num_segments, d_begin_offsets, d_end_offsets,

196  begin_bit, end_bit,

197  stream, debug_synchronous

198  );

199  detail::update_double_buffer(d_keys, d_keys_db);

200  return error;

201  }

202 

203  template<typename KeyT, typename OffsetIteratorT>

204  HIPCUB_RUNTIME_FUNCTION static

205  hipError_t SortKeysDescending(void * d_temp_storage,

206  size_t& temp_storage_bytes,

207  const KeyT * d_keys_in,

208  KeyT * d_keys_out,

209  int num_items,

210  int num_segments,

211  OffsetIteratorT d_begin_offsets,

212  OffsetIteratorT d_end_offsets,

213  int begin_bit = 0,

214  int end_bit = sizeof(KeyT) * 8,

215  hipStream_t stream = 0,

216  bool debug_synchronous = false)

217  {

218  return ::rocprim::segmented_radix_sort_keys_desc(

219  d_temp_storage, temp_storage_bytes,

220  d_keys_in, d_keys_out, num_items,

221  num_segments, d_begin_offsets, d_end_offsets,

222  begin_bit, end_bit,

223  stream, debug_synchronous

224  );

225  }

226 

227  template<typename KeyT, typename OffsetIteratorT>

228  HIPCUB_RUNTIME_FUNCTION static

229  hipError_t SortKeysDescending(void * d_temp_storage,

230  size_t& temp_storage_bytes,

231  DoubleBuffer<KeyT>& d_keys,

232  int num_items,

233  int num_segments,

234  OffsetIteratorT d_begin_offsets,

235  OffsetIteratorT d_end_offsets,

236  int begin_bit = 0,

237  int end_bit = sizeof(KeyT) * 8,

238  hipStream_t stream = 0,

239  bool debug_synchronous = false)

240  {

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

242  hipError_t error = ::rocprim::segmented_radix_sort_keys_desc(

243  d_temp_storage, temp_storage_bytes,

244  d_keys_db, num_items,

245  num_segments, d_begin_offsets, d_end_offsets,

246  begin_bit, end_bit,

247  stream, debug_synchronous

248  );

249  detail::update_double_buffer(d_keys, d_keys_db);

250  return error;

251  }

252 };

253 

254 END_HIPCUB_NAMESPACE

255 

256 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SEGMENTED_RADIX_SORT_HPP_