/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.1/hipcub/include/hipcub/backend/rocprim/block/radix_rank_sort_operations.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/radix_rank_sort_operations.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.5.1/hipcub/include/hipcub/backend/rocprim/block/radix_rank_sort_operations.hpp Source File
radix_rank_sort_operations.hpp
Go to the documentation of this file.
1 /******************************************************************************
2  * Copyright (c) 2011-2020, NVIDIA CORPORATION. All rights reserved.
3  * Modifications Copyright (c) 2021, Advanced Micro Devices, Inc. All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  * * Redistributions of source code must retain the above copyright
8  * notice, this list of conditions and the following disclaimer.
9  * * Redistributions in binary form must reproduce the above copyright
10  * notice, this list of conditions and the following disclaimer in the
11  * documentation and/or other materials provided with the distribution.
12  * * Neither the name of the NVIDIA CORPORATION nor the
13  * names of its contributors may be used to endorse or promote products
14  * derived from this software without specific prior written permission.
15  *
16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26  *
27  ******************************************************************************/
28 
35  #ifndef HIPCUB_ROCPRIM_BLOCK_RADIX_RANK_SORT_OPERATIONS_HPP_
36  #define HIPCUB_ROCPRIM_BLOCK_RADIX_RANK_SORT_OPERATIONS_HPP_
37 
38 #include <type_traits>
39 
40 #include "../../../config.hpp"
41 
42  #include <rocprim/config.hpp>
43  #include <rocprim/type_traits.hpp>
44  #include <rocprim/detail/various.hpp>
45 
46 BEGIN_HIPCUB_NAMESPACE
47 
49 template <bool IS_DESCENDING, typename KeyT>
50 struct RadixSortTwiddle
51 {
52  typedef Traits<KeyT> TraitsT;
53  typedef typename TraitsT::UnsignedBits UnsignedBits;
54  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits In(UnsignedBits key)
55  {
56  key = TraitsT::TwiddleIn(key);
57  if (IS_DESCENDING) key = ~key;
58  return key;
59  }
60  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits Out(UnsignedBits key)
61  {
62  if (IS_DESCENDING) key = ~key;
63  key = TraitsT::TwiddleOut(key);
64  return key;
65  }
66  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits DefaultKey()
67  {
68  return Out(~UnsignedBits(0));
69  }
70 };
71 
88  template <typename KeyT>
89  struct BaseDigitExtractor
90  {
91  typedef Traits<KeyT> TraitsT;
92  typedef typename TraitsT::UnsignedBits UnsignedBits;
93 
94  enum
95  {
96  FLOAT_KEY = TraitsT::CATEGORY == FLOATING_POINT,
97  };
98 
99  static __device__ __forceinline__ UnsignedBits ProcessFloatMinusZero(UnsignedBits key)
100  {
101  if (!FLOAT_KEY) {
102  return key;
103  } else {
104  UnsignedBits TWIDDLED_MINUS_ZERO_BITS =
105  TraitsT::TwiddleIn(UnsignedBits(1) << UnsignedBits(8 * sizeof(UnsignedBits) - 1));
106  UnsignedBits TWIDDLED_ZERO_BITS = TraitsT::TwiddleIn(0);
107  return key == TWIDDLED_MINUS_ZERO_BITS ? TWIDDLED_ZERO_BITS : key;
108  }
109  }
110  };
111 
114  template <typename KeyT>
115  struct BFEDigitExtractor : BaseDigitExtractor<KeyT>
116  {
117  using typename BaseDigitExtractor<KeyT>::UnsignedBits;
118 
119  uint32_t bit_start, num_bits;
120  explicit __device__ __forceinline__ BFEDigitExtractor(
121  uint32_t bit_start = 0, uint32_t num_bits = 0)
122  : bit_start(bit_start), num_bits(num_bits)
123  { }
124 
125  __device__ __forceinline__ uint32_t Digit(UnsignedBits key)
126  {
127  return BFE(this->ProcessFloatMinusZero(key), bit_start, num_bits);
128  }
129  };
130 
133  template <typename KeyT>
134  struct ShiftDigitExtractor : BaseDigitExtractor<KeyT>
135  {
136  using typename BaseDigitExtractor<KeyT>::UnsignedBits;
137 
138  uint32_t bit_start, mask;
139  explicit __device__ __forceinline__ ShiftDigitExtractor(
140  uint32_t bit_start = 0, uint32_t num_bits = 0)
141  : bit_start(bit_start), mask((1 << num_bits) - 1)
142  { }
143 
144  __device__ __forceinline__ uint32_t Digit(UnsignedBits key)
145  {
146  return uint32_t(this->ProcessFloatMinusZero(key) >> UnsignedBits(bit_start)) & mask;
147  }
148  };
149 
150 END_HIPCUB_NAMESPACE
151 
152 #endif //HIPCUB_ROCPRIM_BLOCK_RADIX_RANK_SORT_OPERATIONS_HPP_
hipcub::BFEDigitExtractor
A wrapper type to extract digits. Uses the BFE intrinsic to extract a key from a digit.
Definition: radix_rank_sort_operations.hpp:116
hipcub::BaseDigitExtractor
Base struct for digit extractor. Contains common code to provide special handling for floating-point ...
Definition: radix_rank_sort_operations.hpp:90
hipcub::RadixSortTwiddle
Twiddling keys for radix sort.
Definition: radix_rank_sort_operations.hpp:51
hipcub::ShiftDigitExtractor
A wrapper type to extract digits. Uses a combination of shift and bitwise and to extract digits.
Definition: radix_rank_sort_operations.hpp:135