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

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.6.0/hipcub/include/hipcub/backend/rocprim/block/block_load_func.hpp Source File
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_