/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.7.0/hipcub/include/hipcub/backend/rocprim/util_type.hpp Source File

• home
• docs
• checkouts
• readthedocs.org
• user_builds
• advanced-micro-devices-hipcub
• checkouts
• docs-5.7.0
• hipcub
• include
• hipcub
• backend
• rocprim

util_type.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) 2021, 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_UTIL_TYPE_HPP_

31 #define HIPCUB_ROCPRIM_UTIL_TYPE_HPP_

32 

33 #include <limits>

34 #include <type_traits>

35 

36 #include "../../config.hpp"

37 

38 #include <rocprim/detail/various.hpp>

39 #include <rocprim/types/future_value.hpp>

40 

41 #include <hip/hip_fp16.h>

42 #include <hip/hip_bfloat16.h>

43 

44 BEGIN_HIPCUB_NAMESPACE

45 

46 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

47 

48 using NullType = ::rocprim::empty_type;

49 

50 #endif

51 

52 template<bool B, typename T, typename F> struct

53 [[deprecated("[Since 1.16] If is deprecated use std::conditional instead.")]] If

54 {

55  using Type = typename std::conditional<B, T, F>::type;

56 };

57 

58 template<typename T> struct

59 [[deprecated("[Since 1.16] IsPointer is deprecated use std::is_pointer instead.")]] IsPointer

60 {

61  static constexpr bool VALUE = std::is_pointer<T>::value;

62 };

63 

64 template<typename T> struct

65 [[deprecated("[Since 1.16] IsVolatile is deprecated use std::is_volatile instead.")]] IsVolatile

66 {

67  static constexpr bool VALUE = std::is_volatile<T>::value;

68 };

69 

70 template<typename T> struct

71 [[deprecated("[Since 1.16] RemoveQualifiers is deprecated use std::remove_cv instead.")]] RemoveQualifiers

72 {

73  using Type = typename std::remove_cv<T>::type;

74 };

75 

76 template<int N>

77 struct PowerOfTwo

78 {

79  static constexpr bool VALUE = ::rocprim::detail::is_power_of_two(N);

80 };

81 

82 namespace detail

83 {

84 

85 template<int N, int CURRENT_VAL = N, int COUNT = 0>

86 struct Log2Impl

87 {

88  static constexpr int VALUE = Log2Impl<N, (CURRENT_VAL >> 1), COUNT + 1>::VALUE;

89 };

90 

91 template<int N, int COUNT>

92 struct Log2Impl<N, 0, COUNT>

93 {

94  static constexpr int VALUE = (1 << (COUNT - 1) < N) ? COUNT : COUNT - 1;

95 };

96 

97 } // end of detail namespace

98 

99 template<int N>

100 struct Log2

101 {

102  static_assert(N != 0, "The logarithm of zero is undefined");

103  static constexpr int VALUE = detail::Log2Impl<N>::VALUE;

104 };

105 

106 template<typename T>

107 struct DoubleBuffer

108 {

109  T * d_buffers[2];

110 

111  int selector;

112 

113  HIPCUB_HOST_DEVICE inline

114  DoubleBuffer()

115  {

116  selector = 0;

117  d_buffers[0] = nullptr;

118  d_buffers[1] = nullptr;

119  }

120 

121  HIPCUB_HOST_DEVICE inline

122  DoubleBuffer(T * d_current, T * d_alternate)

123  {

124  selector = 0;

125  d_buffers[0] = d_current;

126  d_buffers[1] = d_alternate;

127  }

128 

129  HIPCUB_HOST_DEVICE inline

130  T * Current()

131  {

132  return d_buffers[selector];

133  }

134 

135  HIPCUB_HOST_DEVICE inline

136  T * Alternate()

137  {

138  return d_buffers[selector ^ 1];

139  }

140 };

141 

142 template <int A>

143 struct Int2Type

144 {

145  enum {VALUE = A};

146 };

147 

148 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

149 

150 template<

151  class Key,

152  class Value

153 >

154 using KeyValuePair = ::rocprim::key_value_pair<Key, Value>;

155 

156 #endif

157 

158 template <typename T, typename Iter = T*>

159 using FutureValue = ::rocprim::future_value<T, Iter>;

160 

161 namespace detail

162 {

163 

164 template<typename T>

165 inline

166 ::rocprim::double_buffer<T> to_double_buffer(DoubleBuffer<T>& source)

167 {

168  return ::rocprim::double_buffer<T>(source.Current(), source.Alternate());

169 }

170 

171 template<typename T>

172 inline

173 void update_double_buffer(DoubleBuffer<T>& target, ::rocprim::double_buffer<T>& source)

174 {

175  if(target.Current() != source.current())

176  {

177  target.selector ^= 1;

178  }

179 }

180 

181 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

182 

183 template <typename T>

184 using is_integral_or_enum =

185  std::integral_constant<bool, std::is_integral<T>::value || std::is_enum<T>::value>;

186 

187 #endif

188 

189 }

190 

191 template <typename NumeratorT, typename DenominatorT>

192 HIPCUB_HOST_DEVICE __forceinline__ constexpr NumeratorT

193 DivideAndRoundUp(NumeratorT n, DenominatorT d)

194 {

195  static_assert(hipcub::detail::is_integral_or_enum<NumeratorT>::value &&

196  hipcub::detail::is_integral_or_enum<DenominatorT>::value,

197  "DivideAndRoundUp is only intended for integral types.");

198 

199  // Static cast to undo integral promotion.

200  return static_cast<NumeratorT>(n / d + (n % d != 0 ? 1 : 0));

201 }

202 

203 #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

204 

205 /******************************************************************************

206  * Size and alignment

207  ******************************************************************************/

208 

210 template <typename T>

211 struct AlignBytes

212 {

213  struct Pad

214  {

215  T val;

216  char byte;

217  };

218 

219  enum

220  {

222  ALIGN_BYTES = sizeof(Pad) - sizeof(T)

223  };

224 

226  typedef T Type;

227 };

228 

229 // Specializations where host C++ compilers (e.g., 32-bit Windows) may disagree

230 // with device C++ compilers (EDG) on types passed as template parameters through

231 // kernel functions

232 

233 #define __HIPCUB_ALIGN_BYTES(t, b) \

234  template <> struct AlignBytes<t> \

235  { enum { ALIGN_BYTES = b }; typedef __align__(b) t Type; };

236 

237 __HIPCUB_ALIGN_BYTES(short4, 8)

238 __HIPCUB_ALIGN_BYTES(ushort4, 8)

239 __HIPCUB_ALIGN_BYTES(int2, 8)

240 __HIPCUB_ALIGN_BYTES(uint2, 8)

241 __HIPCUB_ALIGN_BYTES(long long, 8)

242 __HIPCUB_ALIGN_BYTES(unsigned long long, 8)

243 __HIPCUB_ALIGN_BYTES(float2, 8)

244 __HIPCUB_ALIGN_BYTES(double, 8)

245 #ifdef _WIN32

246  __HIPCUB_ALIGN_BYTES(long2, 8)

247  __HIPCUB_ALIGN_BYTES(ulong2, 8)

248 #else

249  __HIPCUB_ALIGN_BYTES(long2, 16)

250  __HIPCUB_ALIGN_BYTES(ulong2, 16)

251 #endif

252 __HIPCUB_ALIGN_BYTES(int4, 16)

253 __HIPCUB_ALIGN_BYTES(uint4, 16)

254 __HIPCUB_ALIGN_BYTES(float4, 16)

255 __HIPCUB_ALIGN_BYTES(long4, 16)

256 __HIPCUB_ALIGN_BYTES(ulong4, 16)

257 __HIPCUB_ALIGN_BYTES(longlong2, 16)

258 __HIPCUB_ALIGN_BYTES(ulonglong2, 16)

259 __HIPCUB_ALIGN_BYTES(double2, 16)

260 __HIPCUB_ALIGN_BYTES(longlong4, 16)

261 __HIPCUB_ALIGN_BYTES(ulonglong4, 16)

262 __HIPCUB_ALIGN_BYTES(double4, 16)

263 

264 template <typename T> struct AlignBytes<volatile T> : AlignBytes<T> {};

265 template <typename T> struct AlignBytes<const T> : AlignBytes<T> {};

266 template <typename T> struct AlignBytes<const volatile T> : AlignBytes<T> {};

267 

268 

270 template <typename T>

271 struct UnitWord

272 {

273  enum {

274  ALIGN_BYTES = AlignBytes<T>::ALIGN_BYTES

275  };

276 

277  template <typename Unit>

278  struct IsMultiple

279  {

280  enum {

281  UNIT_ALIGN_BYTES = AlignBytes<Unit>::ALIGN_BYTES,

282  IS_MULTIPLE = (sizeof(T) % sizeof(Unit) == 0) && (int(ALIGN_BYTES) % int(UNIT_ALIGN_BYTES) == 0)

283  };

284  };

285 

287  typedef typename std::conditional<IsMultiple<int>::IS_MULTIPLE,

288  unsigned int,

289  typename std::conditional<IsMultiple<short>::IS_MULTIPLE,

290  unsigned short,

291  unsigned char>::type>::type ShuffleWord;

292 

294  typedef typename std::conditional<IsMultiple<long long>::IS_MULTIPLE,

295  unsigned long long,

296  ShuffleWord>::type VolatileWord;

297 

299  typedef typename std::conditional<IsMultiple<longlong2>::IS_MULTIPLE,

300  ulonglong2,

301  VolatileWord>::type DeviceWord;

302 

304  typedef typename std::conditional<IsMultiple<int4>::IS_MULTIPLE,

305  uint4,

306  typename std::conditional<IsMultiple<int2>::IS_MULTIPLE,

307  uint2,

308  ShuffleWord>::type>::type TextureWord;

309 };

310 

311 

312 // float2 specialization workaround (for SM10-SM13)

313 template <>

314 struct UnitWord <float2>

315 {

316  typedef int ShuffleWord;

317  typedef unsigned long long VolatileWord;

318  typedef unsigned long long DeviceWord;

319  typedef float2 TextureWord;

320 };

321 

322 // float4 specialization workaround (for SM10-SM13)

323 template <>

324 struct UnitWord <float4>

325 {

326  typedef int ShuffleWord;

327  typedef unsigned long long VolatileWord;

328  typedef ulonglong2 DeviceWord;

329  typedef float4 TextureWord;

330 };

331 

332 

333 // char2 specialization workaround (for SM10-SM13)

334 template <>

335 struct UnitWord <char2>

336 {

337  typedef unsigned short ShuffleWord;

338  typedef unsigned short VolatileWord;

339  typedef unsigned short DeviceWord;

340  typedef unsigned short TextureWord;

341 };

342 

343 

344 template <typename T> struct UnitWord<volatile T> : UnitWord<T> {};

345 template <typename T> struct UnitWord<const T> : UnitWord<T> {};

346 template <typename T> struct UnitWord<const volatile T> : UnitWord<T> {};

347 

348 

349 #endif // DOXYGEN_SHOULD_SKIP_THIS

350 

351 

352 

353 

354 /******************************************************************************

355  * Wrapper types

356  ******************************************************************************/

357 

361 template <typename T>

362 struct Uninitialized

363 {

365  typedef typename UnitWord<T>::DeviceWord DeviceWord;

366 

367  static constexpr std::size_t DATA_SIZE = sizeof(T);

368  static constexpr std::size_t WORD_SIZE = sizeof(DeviceWord);

369  static constexpr std::size_t WORDS = DATA_SIZE / WORD_SIZE;

370 

372  DeviceWord storage[WORDS];

373 

375  HIPCUB_HOST_DEVICE __forceinline__ T& Alias()

376  {

377  return reinterpret_cast<T&>(*this);

378  }

379 };

380 

381 

382 /******************************************************************************

383  * Simple type traits utilities.

384  *

385  * For example:

386  * Traits<int>::CATEGORY // SIGNED_INTEGER

387  * Traits<NullType>::NULL_TYPE // true

388  * Traits<uint4>::CATEGORY // NOT_A_NUMBER

389  * Traits<uint4>::PRIMITIVE; // false

390  *

391  ******************************************************************************/

392 

393  #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document

394 

398 enum Category

399 {

400  NOT_A_NUMBER,

401  SIGNED_INTEGER,

402  UNSIGNED_INTEGER,

403  FLOATING_POINT

404 };

405 

406 

410 template <Category _CATEGORY, bool _PRIMITIVE, bool _NULL_TYPE, typename _UnsignedBits, typename T>

411 struct BaseTraits

412 {

414  static const Category CATEGORY = _CATEGORY;

415  enum

416  {

417  PRIMITIVE = _PRIMITIVE,

418  NULL_TYPE = _NULL_TYPE,

419  };

420 };

421 

422 

426 template <typename _UnsignedBits, typename T>

427 struct BaseTraits<UNSIGNED_INTEGER, true, false, _UnsignedBits, T>

428 {

429  typedef _UnsignedBits UnsignedBits;

430 

431  static const Category CATEGORY = UNSIGNED_INTEGER;

432  static const UnsignedBits LOWEST_KEY = UnsignedBits(0);

433  static const UnsignedBits MAX_KEY = UnsignedBits(-1);

434 

435  enum

436  {

437  PRIMITIVE = true,

438  NULL_TYPE = false,

439  };

440 

441 

442  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleIn(UnsignedBits key)

443  {

444  return key;

445  }

446 

447  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleOut(UnsignedBits key)

448  {

449  return key;

450  }

451 

452  static HIPCUB_HOST_DEVICE __forceinline__ T Max()

453  {

454  UnsignedBits retval_bits = MAX_KEY;

455  T retval;

456  memcpy(&retval, &retval_bits, sizeof(T));

457  return retval;

458  }

459 

460  static HIPCUB_HOST_DEVICE __forceinline__ T Lowest()

461  {

462  UnsignedBits retval_bits = LOWEST_KEY;

463  T retval;

464  memcpy(&retval, &retval_bits, sizeof(T));

465  return retval;

466  }

467 };

468 

469 

473 template <typename _UnsignedBits, typename T>

474 struct BaseTraits<SIGNED_INTEGER, true, false, _UnsignedBits, T>

475 {

476  typedef _UnsignedBits UnsignedBits;

477 

478  static const Category CATEGORY = SIGNED_INTEGER;

479  static const UnsignedBits HIGH_BIT = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);

480  static const UnsignedBits LOWEST_KEY = HIGH_BIT;

481  static const UnsignedBits MAX_KEY = UnsignedBits(-1) ^ HIGH_BIT;

482 

483  enum

484  {

485  PRIMITIVE = true,

486  NULL_TYPE = false,

487  };

488 

489  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleIn(UnsignedBits key)

490  {

491  return key ^ HIGH_BIT;

492  };

493 

494  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleOut(UnsignedBits key)

495  {

496  return key ^ HIGH_BIT;

497  };

498 

499  static HIPCUB_HOST_DEVICE __forceinline__ T Max()

500  {

501  UnsignedBits retval = MAX_KEY;

502  return reinterpret_cast<T&>(retval);

503  }

504 

505  static HIPCUB_HOST_DEVICE __forceinline__ T Lowest()

506  {

507  UnsignedBits retval = LOWEST_KEY;

508  return reinterpret_cast<T&>(retval);

509  }

510 };

511 

512 template <typename _T>

513 struct FpLimits;

514 

515 template <>

516 struct FpLimits<float>

517 {

518  static HIPCUB_HOST_DEVICE __forceinline__ float Max() {

519  return std::numeric_limits<float>::max();

520  }

521 

522  static HIPCUB_HOST_DEVICE __forceinline__ float Lowest() {

523  return std::numeric_limits<float>::max() * float(-1);

524  }

525 };

526 

527 template <>

528 struct FpLimits<double>

529 {

530  static HIPCUB_HOST_DEVICE __forceinline__ double Max() {

531  return std::numeric_limits<double>::max();

532  }

533 

534  static HIPCUB_HOST_DEVICE __forceinline__ double Lowest() {

535  return std::numeric_limits<double>::max() * double(-1);

536  }

537 };

538 

539 template <>

540 struct FpLimits<__half>

541 {

542  static HIPCUB_HOST_DEVICE __forceinline__ __half Max() {

543  unsigned short max_word = 0x7BFF;

544  return reinterpret_cast<__half&>(max_word);

545  }

546 

547  static HIPCUB_HOST_DEVICE __forceinline__ __half Lowest() {

548  unsigned short lowest_word = 0xFBFF;

549  return reinterpret_cast<__half&>(lowest_word);

550  }

551 };

552 

553 template <>

554 struct FpLimits<hip_bfloat16>

555 {

556  static HIPCUB_HOST_DEVICE __forceinline__ hip_bfloat16 Max() {

557  unsigned short max_word = 0x7F7F;

558  return reinterpret_cast<hip_bfloat16 &>(max_word);

559  }

560 

561  static HIPCUB_HOST_DEVICE __forceinline__ hip_bfloat16 Lowest() {

562  unsigned short lowest_word = 0xFF7F;

563  return reinterpret_cast<hip_bfloat16 &>(lowest_word);

564  }

565 };

566 

570 template <typename _UnsignedBits, typename T>

571 struct BaseTraits<FLOATING_POINT, true, false, _UnsignedBits, T>

572 {

573  typedef _UnsignedBits UnsignedBits;

574 

575  static const Category CATEGORY = FLOATING_POINT;

576  static const UnsignedBits HIGH_BIT = UnsignedBits(1) << ((sizeof(UnsignedBits) * 8) - 1);

577  static const UnsignedBits LOWEST_KEY = UnsignedBits(-1);

578  static const UnsignedBits MAX_KEY = UnsignedBits(-1) ^ HIGH_BIT;

579 

580  enum

581  {

582  PRIMITIVE = true,

583  NULL_TYPE = false,

584  };

585 

586  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleIn(UnsignedBits key)

587  {

588  UnsignedBits mask = (key & HIGH_BIT) ? UnsignedBits(-1) : HIGH_BIT;

589  return key ^ mask;

590  };

591 

592  static HIPCUB_HOST_DEVICE __forceinline__ UnsignedBits TwiddleOut(UnsignedBits key)

593  {

594  UnsignedBits mask = (key & HIGH_BIT) ? HIGH_BIT : UnsignedBits(-1);

595  return key ^ mask;

596  };

597 

598  static HIPCUB_HOST_DEVICE __forceinline__ T Max() {

599  return FpLimits<T>::Max();

600  }

601 

602  static HIPCUB_HOST_DEVICE __forceinline__ T Lowest() {

603  return FpLimits<T>::Lowest();

604  }

605 };

606 

607 

611 template <typename T> struct NumericTraits : BaseTraits<NOT_A_NUMBER, false, false, T, T> {};

612 

613 template <> struct NumericTraits<NullType> : BaseTraits<NOT_A_NUMBER, false, true, NullType, NullType> {};

614 

615 template <> struct NumericTraits<char> : BaseTraits<(std::numeric_limits<char>::is_signed) ? SIGNED_INTEGER : UNSIGNED_INTEGER, true, false, unsigned char, char> {};

616 template <> struct NumericTraits<signed char> : BaseTraits<SIGNED_INTEGER, true, false, unsigned char, signed char> {};

617 template <> struct NumericTraits<short> : BaseTraits<SIGNED_INTEGER, true, false, unsigned short, short> {};

618 template <> struct NumericTraits<int> : BaseTraits<SIGNED_INTEGER, true, false, unsigned int, int> {};

619 template <> struct NumericTraits<long> : BaseTraits<SIGNED_INTEGER, true, false, unsigned long, long> {};

620 template <> struct NumericTraits<long long> : BaseTraits<SIGNED_INTEGER, true, false, unsigned long long, long long> {};

621 

622 template <> struct NumericTraits<unsigned char> : BaseTraits<UNSIGNED_INTEGER, true, false, unsigned char, unsigned char> {};

623 template <> struct NumericTraits<unsigned short> : BaseTraits<UNSIGNED_INTEGER, true, false, unsigned short, unsigned short> {};

624 template <> struct NumericTraits<unsigned int> : BaseTraits<UNSIGNED_INTEGER, true, false, unsigned int, unsigned int> {};

625 template <> struct NumericTraits<unsigned long> : BaseTraits<UNSIGNED_INTEGER, true, false, unsigned long, unsigned long> {};

626 template <> struct NumericTraits<unsigned long long> : BaseTraits<UNSIGNED_INTEGER, true, false, unsigned long long, unsigned long long> {};

627 

628 template <> struct NumericTraits<float> : BaseTraits<FLOATING_POINT, true, false, unsigned int, float> {};

629 template <> struct NumericTraits<double> : BaseTraits<FLOATING_POINT, true, false, unsigned long long, double> {};

630 template <> struct NumericTraits<__half> : BaseTraits<FLOATING_POINT, true, false, unsigned short, __half> {};

631 template <> struct NumericTraits<hip_bfloat16 > : BaseTraits<FLOATING_POINT, true, false, unsigned short, hip_bfloat16 > {};

632 

633 template <> struct NumericTraits<bool> : BaseTraits<UNSIGNED_INTEGER, true, false, typename UnitWord<bool>::VolatileWord, bool> {};

634 

638 template <typename T>

639 struct Traits : NumericTraits<typename std::remove_cv<T>::type> {};

640 

641 #endif // DOXYGEN_SHOULD_SKIP_THIS

642 

643 END_HIPCUB_NAMESPACE

644 

645 #endif // HIPCUB_ROCPRIM_UTIL_TYPE_HPP_

hipcub::DoubleBuffer

Definition: util_type.hpp:108

hipcub::If

Definition: util_type.hpp:54

hipcub::Int2Type

Definition: util_type.hpp:144

hipcub::IsPointer

Definition: util_type.hpp:60

hipcub::IsVolatile

Definition: util_type.hpp:66

hipcub::Log2

Definition: util_type.hpp:101

hipcub::PowerOfTwo

Definition: util_type.hpp:78

hipcub::RemoveQualifiers

Definition: util_type.hpp:72

hipcub::Uninitialized

A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions.

Definition: util_type.hpp:363

hipcub::Uninitialized::Alias

__host__ __device__ __forceinline__ T & Alias()

Alias.

Definition: util_type.hpp:375

hipcub::Uninitialized::DeviceWord

UnitWord< T >::DeviceWord DeviceWord

Biggest memory-access word that T is a whole multiple of and is not larger than the alignment of T.

Definition: util_type.hpp:365