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

block_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_BLOCK_BLOCK_RADIX_SORT_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_RADIX_SORT_HPP_

32 

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

34 

35 #include "../util_type.hpp"

36 

37 #include <rocprim/functional.hpp>

38 #include <rocprim/block/block_radix_sort.hpp>

39 

40 #include "block_scan.hpp"

41 

42 BEGIN_HIPCUB_NAMESPACE

43 

44 template<

45  typename KeyT,

46  int BLOCK_DIM_X,

47  int ITEMS_PER_THREAD,

48  typename ValueT = NullType,

49  int RADIX_BITS = 4, /* ignored */

50  bool MEMOIZE_OUTER_SCAN = true, /* ignored */

51  BlockScanAlgorithm INNER_SCAN_ALGORITHM = BLOCK_SCAN_WARP_SCANS, /* ignored */

52  hipSharedMemConfig SMEM_CONFIG = hipSharedMemBankSizeFourByte, /* ignored */

53  int BLOCK_DIM_Y = 1,

54  int BLOCK_DIM_Z = 1,

55  int PTX_ARCH = HIPCUB_ARCH /* ignored */

56 >

57 class BlockRadixSort

58  : private ::rocprim::block_radix_sort<

59  KeyT,

60  BLOCK_DIM_X,

61  ITEMS_PER_THREAD,

62  ValueT,

63  BLOCK_DIM_Y,

64  BLOCK_DIM_Z

65  >

66 {

67  static_assert(

68  BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z > 0,

69  "BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z must be greater than 0"

70  );

71 

72  using base_type =

73  typename ::rocprim::block_radix_sort<

74  KeyT,

75  BLOCK_DIM_X,

76  ITEMS_PER_THREAD,

77  ValueT,

78  BLOCK_DIM_Y,

79  BLOCK_DIM_Z

80  >;

81 

82  // Reference to temporary storage (usually shared memory)

83  typename base_type::storage_type& temp_storage_;

84 

85 public:

86  using TempStorage = typename base_type::storage_type;

87 

88  HIPCUB_DEVICE inline

89  BlockRadixSort() : temp_storage_(private_storage())

90  {

91  }

92 

93  HIPCUB_DEVICE inline

94  BlockRadixSort(TempStorage& temp_storage) : temp_storage_(temp_storage)

95  {

96  }

97 

98  HIPCUB_DEVICE inline

99  void Sort(KeyT (&keys)[ITEMS_PER_THREAD],

100  int begin_bit = 0,

101  int end_bit = sizeof(KeyT) * 8)

102  {

103  base_type::sort(keys, temp_storage_, begin_bit, end_bit);

104  }

105 

106  HIPCUB_DEVICE inline

107  void Sort(KeyT (&keys)[ITEMS_PER_THREAD],

108  ValueT (&values)[ITEMS_PER_THREAD],

109  int begin_bit = 0,

110  int end_bit = sizeof(KeyT) * 8)

111  {

112  base_type::sort(keys, values, temp_storage_, begin_bit, end_bit);

113  }

114 

115  HIPCUB_DEVICE inline

116  void SortDescending(KeyT (&keys)[ITEMS_PER_THREAD],

117  int begin_bit = 0,

118  int end_bit = sizeof(KeyT) * 8)

119  {

120  base_type::sort_desc(keys, temp_storage_, begin_bit, end_bit);

121  }

122 

123  HIPCUB_DEVICE inline

124  void SortDescending(KeyT (&keys)[ITEMS_PER_THREAD],

125  ValueT (&values)[ITEMS_PER_THREAD],

126  int begin_bit = 0,

127  int end_bit = sizeof(KeyT) * 8)

128  {

129  base_type::sort_desc(keys, values, temp_storage_, begin_bit, end_bit);

130  }

131 

132  HIPCUB_DEVICE inline

133  void SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],

134  int begin_bit = 0,

135  int end_bit = sizeof(KeyT) * 8)

136  {

137  base_type::sort_to_striped(keys, temp_storage_, begin_bit, end_bit);

138  }

139 

140  HIPCUB_DEVICE inline

141  void SortBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],

142  ValueT (&values)[ITEMS_PER_THREAD],

143  int begin_bit = 0,

144  int end_bit = sizeof(KeyT) * 8)

145  {

146  base_type::sort_to_striped(keys, values, temp_storage_, begin_bit, end_bit);

147  }

148 

149  HIPCUB_DEVICE inline

150  void SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],

151  int begin_bit = 0,

152  int end_bit = sizeof(KeyT) * 8)

153  {

154  base_type::sort_desc_to_striped(keys, temp_storage_, begin_bit, end_bit);

155  }

156 

157  HIPCUB_DEVICE inline

158  void SortDescendingBlockedToStriped(KeyT (&keys)[ITEMS_PER_THREAD],

159  ValueT (&values)[ITEMS_PER_THREAD],

160  int begin_bit = 0,

161  int end_bit = sizeof(KeyT) * 8)

162  {

163  base_type::sort_desc_to_striped(keys, values, temp_storage_, begin_bit, end_bit);

164  }

165 

166 private:

167  HIPCUB_DEVICE inline

168  TempStorage& private_storage()

169  {

170  HIPCUB_SHARED_MEMORY TempStorage private_storage;

171  return private_storage;

172  }

173 };

174 

175 END_HIPCUB_NAMESPACE

176 

177 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_RADIX_SORT_HPP_

hipcub::BlockRadixSort

Definition: block_radix_sort.hpp:66