/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/block/block_store.hpp Source File

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

block_store.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_STORE_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_STORE_HPP_

32 

33 #include <type_traits>

34 

35 #include "../../../config.hpp"

36 

37 #include "block_store_func.hpp"

38 

39 #include <rocprim/block/block_store.hpp>

40 

41 BEGIN_HIPCUB_NAMESPACE

42 

43 namespace detail

44 {

45  inline constexpr

46  typename std::underlying_type<::rocprim::block_store_method>::type

47  to_BlockStoreAlgorithm_enum(::rocprim::block_store_method v)

48  {

49  using utype = std::underlying_type<::rocprim::block_store_method>::type;

50  return static_cast<utype>(v);

51  }

52 }

53 

54 enum BlockStoreAlgorithm

55 {

56  BLOCK_STORE_DIRECT

57  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_direct),

58  BLOCK_STORE_STRIPED

59  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_striped),

60  BLOCK_STORE_VECTORIZE

61  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_vectorize),

62  BLOCK_STORE_TRANSPOSE

63  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_transpose),

64  BLOCK_STORE_WARP_TRANSPOSE

65  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_warp_transpose),

66  BLOCK_STORE_WARP_TRANSPOSE_TIMESLICED

67  = detail::to_BlockStoreAlgorithm_enum(::rocprim::block_store_method::block_store_warp_transpose)

68 };

69 

70 template<

71  typename T,

72  int BLOCK_DIM_X,

73  int ITEMS_PER_THREAD,

74  BlockStoreAlgorithm ALGORITHM = BLOCK_STORE_DIRECT,

75  int BLOCK_DIM_Y = 1,

76  int BLOCK_DIM_Z = 1,

77  int ARCH = HIPCUB_ARCH /* ignored */

78 >

79 class BlockStore

80  : private ::rocprim::block_store<

81  T,

82  BLOCK_DIM_X,

83  ITEMS_PER_THREAD,

84  static_cast<::rocprim::block_store_method>(ALGORITHM),

85  BLOCK_DIM_Y,

86  BLOCK_DIM_Z

87  >

88 {

89  static_assert(

90  BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z > 0,

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

92  );

93 

94  using base_type =

95  typename ::rocprim::block_store<

96  T,

97  BLOCK_DIM_X,

98  ITEMS_PER_THREAD,

99  static_cast<::rocprim::block_store_method>(ALGORITHM),

100  BLOCK_DIM_Y,

101  BLOCK_DIM_Z

102  >;

103 

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

105  typename base_type::storage_type& temp_storage_;

106 

107 public:

108  using TempStorage = typename base_type::storage_type;

109 

110  HIPCUB_DEVICE inline

111  BlockStore() : temp_storage_(private_storage())

112  {

113  }

114 

115  HIPCUB_DEVICE inline

116  BlockStore(TempStorage& temp_storage) : temp_storage_(temp_storage)

117  {

118  }

119 

120  template<class OutputIteratorT>

121  HIPCUB_DEVICE inline

122  void Store(OutputIteratorT block_iter,

123  T (&items)[ITEMS_PER_THREAD])

124  {

125  base_type::store(block_iter, items, temp_storage_);

126  }

127 

128  template<class OutputIteratorT>

129  HIPCUB_DEVICE inline

130  void Store(OutputIteratorT block_iter,

131  T (&items)[ITEMS_PER_THREAD],

132  int valid_items)

133  {

134  base_type::store(block_iter, items, valid_items, temp_storage_);

135  }

136 

137 private:

138  HIPCUB_DEVICE inline

139  TempStorage& private_storage()

140  {

141  HIPCUB_SHARED_MEMORY TempStorage private_storage;

142  return private_storage;

143  }

144 };

145 

146 END_HIPCUB_NAMESPACE

147 

148 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_STORE_HPP_

hipcub::BlockStore

Definition: block_store.hpp:88