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

31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_

32 

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

34 

35 #include "../util_type.hpp"

36 

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

38 

39 BEGIN_HIPCUB_NAMESPACE

40 

41 struct DeviceRadixSort

42 {

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

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

52  int begin_bit = 0,

53  int end_bit = sizeof(KeyT) * 8,

54  hipStream_t stream = 0,

55  bool debug_synchronous = false)

56  {

57  return ::rocprim::radix_sort_pairs(

58  d_temp_storage, temp_storage_bytes,

59  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

60  begin_bit, end_bit,

61  stream, debug_synchronous

62  );

63  }

64 

65  template<typename KeyT, typename ValueT, typename NumItemsT>

66  HIPCUB_RUNTIME_FUNCTION static

67  hipError_t SortPairs(void * d_temp_storage,

68  size_t& temp_storage_bytes,

69  DoubleBuffer<KeyT>& d_keys,

70  DoubleBuffer<ValueT>& d_values,

71  NumItemsT num_items,

72  int begin_bit = 0,

73  int end_bit = sizeof(KeyT) * 8,

74  hipStream_t stream = 0,

75  bool debug_synchronous = false)

76  {

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

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

79  hipError_t error = ::rocprim::radix_sort_pairs(

80  d_temp_storage, temp_storage_bytes,

81  d_keys_db, d_values_db, num_items,

82  begin_bit, end_bit,

83  stream, debug_synchronous

84  );

85  detail::update_double_buffer(d_keys, d_keys_db);

86  detail::update_double_buffer(d_values, d_values_db);

87  return error;

88  }

89 

90  template<typename KeyT, typename ValueT, typename NumItemsT>

91  HIPCUB_RUNTIME_FUNCTION static

92  hipError_t SortPairsDescending(void * d_temp_storage,

93  size_t& temp_storage_bytes,

94  const KeyT * d_keys_in,

95  KeyT * d_keys_out,

96  const ValueT * d_values_in,

97  ValueT * d_values_out,

98  NumItemsT num_items,

99  int begin_bit = 0,

100  int end_bit = sizeof(KeyT) * 8,

101  hipStream_t stream = 0,

102  bool debug_synchronous = false)

103  {

104  return ::rocprim::radix_sort_pairs_desc(

105  d_temp_storage, temp_storage_bytes,

106  d_keys_in, d_keys_out, d_values_in, d_values_out, num_items,

107  begin_bit, end_bit,

108  stream, debug_synchronous

109  );

110  }

111 

112  template<typename KeyT, typename ValueT, typename NumItemsT>

113  HIPCUB_RUNTIME_FUNCTION static

114  hipError_t SortPairsDescending(void * d_temp_storage,

115  size_t& temp_storage_bytes,

116  DoubleBuffer<KeyT>& d_keys,

117  DoubleBuffer<ValueT>& d_values,

118  NumItemsT num_items,

119  int begin_bit = 0,

120  int end_bit = sizeof(KeyT) * 8,

121  hipStream_t stream = 0,

122  bool debug_synchronous = false)

123  {

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

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

126  hipError_t error = ::rocprim::radix_sort_pairs_desc(

127  d_temp_storage, temp_storage_bytes,

128  d_keys_db, d_values_db, num_items,

129  begin_bit, end_bit,

130  stream, debug_synchronous

131  );

132  detail::update_double_buffer(d_keys, d_keys_db);

133  detail::update_double_buffer(d_values, d_values_db);

134  return error;

135  }

136 

137  template<typename KeyT, typename NumItemsT>

138  HIPCUB_RUNTIME_FUNCTION static

139  hipError_t SortKeys(void * d_temp_storage,

140  size_t& temp_storage_bytes,

141  const KeyT * d_keys_in,

142  KeyT * d_keys_out,

143  NumItemsT num_items,

144  int begin_bit = 0,

145  int end_bit = sizeof(KeyT) * 8,

146  hipStream_t stream = 0,

147  bool debug_synchronous = false)

148  {

149  return ::rocprim::radix_sort_keys(

150  d_temp_storage, temp_storage_bytes,

151  d_keys_in, d_keys_out, num_items,

152  begin_bit, end_bit,

153  stream, debug_synchronous

154  );

155  }

156 

157  template<typename KeyT, typename NumItemsT>

158  HIPCUB_RUNTIME_FUNCTION static

159  hipError_t SortKeys(void * d_temp_storage,

160  size_t& temp_storage_bytes,

161  DoubleBuffer<KeyT>& d_keys,

162  NumItemsT num_items,

163  int begin_bit = 0,

164  int end_bit = sizeof(KeyT) * 8,

165  hipStream_t stream = 0,

166  bool debug_synchronous = false)

167  {

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

169  hipError_t error = ::rocprim::radix_sort_keys(

170  d_temp_storage, temp_storage_bytes,

171  d_keys_db, num_items,

172  begin_bit, end_bit,

173  stream, debug_synchronous

174  );

175  detail::update_double_buffer(d_keys, d_keys_db);

176  return error;

177  }

178 

179  template<typename KeyT, typename NumItemsT>

180  HIPCUB_RUNTIME_FUNCTION static

181  hipError_t SortKeysDescending(void * d_temp_storage,

182  size_t& temp_storage_bytes,

183  const KeyT * d_keys_in,

184  KeyT * d_keys_out,

185  NumItemsT num_items,

186  int begin_bit = 0,

187  int end_bit = sizeof(KeyT) * 8,

188  hipStream_t stream = 0,

189  bool debug_synchronous = false)

190  {

191  return ::rocprim::radix_sort_keys_desc(

192  d_temp_storage, temp_storage_bytes,

193  d_keys_in, d_keys_out, num_items,

194  begin_bit, end_bit,

195  stream, debug_synchronous

196  );

197  }

198 

199  template<typename KeyT, typename NumItemsT>

200  HIPCUB_RUNTIME_FUNCTION static

201  hipError_t SortKeysDescending(void * d_temp_storage,

202  size_t& temp_storage_bytes,

203  DoubleBuffer<KeyT>& d_keys,

204  NumItemsT num_items,

205  int begin_bit = 0,

206  int end_bit = sizeof(KeyT) * 8,

207  hipStream_t stream = 0,

208  bool debug_synchronous = false)

209  {

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

211  hipError_t error = ::rocprim::radix_sort_keys_desc(

212  d_temp_storage, temp_storage_bytes,

213  d_keys_db, num_items,

214  begin_bit, end_bit,

215  stream, debug_synchronous

216  );

217  detail::update_double_buffer(d_keys, d_keys_db);

218  return error;

219  }

220 };

221 

222 END_HIPCUB_NAMESPACE

223 

224 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_RADIX_SORT_HPP_