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

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

block_load_func.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_LOAD_FUNC_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_LOAD_FUNC_HPP_

32 

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

34 

35 #include <rocprim/block/block_load_func.hpp>

36 

37 BEGIN_HIPCUB_NAMESPACE

38 

39 template<

40  typename T,

41  int ITEMS_PER_THREAD,

42  typename InputIteratorT

43 >

44 HIPCUB_DEVICE inline

45 void LoadDirectBlocked(int linear_id,

46  InputIteratorT block_iter,

47  T (&items)[ITEMS_PER_THREAD])

48 {

49  ::rocprim::block_load_direct_blocked(

50  linear_id, block_iter, items

51  );

52 }

53 

54 template<

55  typename T,

56  int ITEMS_PER_THREAD,

57  typename InputIteratorT

58 >

59 HIPCUB_DEVICE inline

60 void LoadDirectBlocked(int linear_id,

61  InputIteratorT block_iter,

62  T (&items)[ITEMS_PER_THREAD],

63  int valid_items)

64 {

65  ::rocprim::block_load_direct_blocked(

66  linear_id, block_iter, items, valid_items

67  );

68 }

69 

70 template<

71  typename T,

72  typename Default,

73  int ITEMS_PER_THREAD,

74  typename InputIteratorT

75 >

76 HIPCUB_DEVICE inline

77 void LoadDirectBlocked(int linear_id,

78  InputIteratorT block_iter,

79  T (&items)[ITEMS_PER_THREAD],

80  int valid_items,

81  Default oob_default)

82 {

83  ::rocprim::block_load_direct_blocked(

84  linear_id, block_iter, items, valid_items, oob_default

85  );

86 }

87 

88 template <

89  typename T,

90  int ITEMS_PER_THREAD

91 >

92 HIPCUB_DEVICE inline

93 void LoadDirectBlockedVectorized(int linear_id,

94  T* block_iter,

95  T (&items)[ITEMS_PER_THREAD])

96 {

97  ::rocprim::block_load_direct_blocked_vectorized(

98  linear_id, block_iter, items

99  );

100 }

101 

102 template<

103  int BLOCK_THREADS,

104  typename T,

105  int ITEMS_PER_THREAD,

106  typename InputIteratorT

107 >

108 HIPCUB_DEVICE inline

109 void LoadDirectStriped(int linear_id,

110  InputIteratorT block_iter,

111  T (&items)[ITEMS_PER_THREAD])

112 {

113  ::rocprim::block_load_direct_striped<BLOCK_THREADS>(

114  linear_id, block_iter, items

115  );

116 }

117 

118 template<

119  int BLOCK_THREADS,

120  typename T,

121  int ITEMS_PER_THREAD,

122  typename InputIteratorT

123 >

124 HIPCUB_DEVICE inline

125 void LoadDirectStriped(int linear_id,

126  InputIteratorT block_iter,

127  T (&items)[ITEMS_PER_THREAD],

128  int valid_items)

129 {

130  ::rocprim::block_load_direct_striped<BLOCK_THREADS>(

131  linear_id, block_iter, items, valid_items

132  );

133 }

134 

135 template<

136  int BLOCK_THREADS,

137  typename T,

138  typename Default,

139  int ITEMS_PER_THREAD,

140  typename InputIteratorT

141 >

142 HIPCUB_DEVICE inline

143 void LoadDirectStriped(int linear_id,

144  InputIteratorT block_iter,

145  T (&items)[ITEMS_PER_THREAD],

146  int valid_items,

147  Default oob_default)

148 {

149  ::rocprim::block_load_direct_striped<BLOCK_THREADS>(

150  linear_id, block_iter, items, valid_items, oob_default

151  );

152 }

153 

154 template<

155  typename T,

156  int ITEMS_PER_THREAD,

157  typename InputIteratorT

158 >

159 HIPCUB_DEVICE inline

160 void LoadDirectWarpStriped(int linear_id,

161  InputIteratorT block_iter,

162  T (&items)[ITEMS_PER_THREAD])

163 {

164  ::rocprim::block_load_direct_warp_striped(

165  linear_id, block_iter, items

166  );

167 }

168 

169 template<

170  typename T,

171  int ITEMS_PER_THREAD,

172  typename InputIteratorT

173 >

174 HIPCUB_DEVICE inline

175 void LoadDirectWarpStriped(int linear_id,

176  InputIteratorT block_iter,

177  T (&items)[ITEMS_PER_THREAD],

178  int valid_items)

179 {

180  ::rocprim::block_load_direct_warp_striped(

181  linear_id, block_iter, items, valid_items

182  );

183 }

184 

185 template<

186  typename T,

187  typename Default,

188  int ITEMS_PER_THREAD,

189  typename InputIteratorT

190 >

191 HIPCUB_DEVICE inline

192 void LoadDirectWarpStriped(int linear_id,

193  InputIteratorT block_iter,

194  T (&items)[ITEMS_PER_THREAD],

195  int valid_items,

196  Default oob_default)

197 {

198  ::rocprim::block_load_direct_warp_striped(

199  linear_id, block_iter, items, valid_items, oob_default

200  );

201 }

202 

203 END_HIPCUB_NAMESPACE

204 

205 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_LOAD_FUNC_HPP_