/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.3.3/hipcub/include/hipcub/backend/rocprim/warp/warp_reduce.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.3.3/hipcub/include/hipcub/backend/rocprim/warp/warp_reduce.hpp Source File#

hipCUB: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.3.3/hipcub/include/hipcub/backend/rocprim/warp/warp_reduce.hpp Source File
warp_reduce.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_WARP_WARP_REDUCE_HPP_
31 #define HIPCUB_ROCPRIM_WARP_WARP_REDUCE_HPP_
32 
33 #include "../../../config.hpp"
34 
35 #include "../util_ptx.hpp"
36 #include "../thread/thread_operators.hpp"
37 
38 #include <rocprim/warp/warp_reduce.hpp>
39 
40 BEGIN_HIPCUB_NAMESPACE
41 
42 template<
43  typename T,
44  int LOGICAL_WARP_THREADS = HIPCUB_DEVICE_WARP_THREADS,
45  int ARCH = HIPCUB_ARCH>
46 class WarpReduce : private ::rocprim::warp_reduce<T, LOGICAL_WARP_THREADS>
47 {
48  static_assert(LOGICAL_WARP_THREADS > 0, "LOGICAL_WARP_THREADS must be greater than 0");
49  using base_type = typename ::rocprim::warp_reduce<T, LOGICAL_WARP_THREADS>;
50 
51  typename base_type::storage_type &temp_storage_;
52 
53 public:
54  using TempStorage = typename base_type::storage_type;
55 
56  HIPCUB_DEVICE inline
57  WarpReduce(TempStorage& temp_storage) : temp_storage_(temp_storage)
58  {
59  }
60 
61  HIPCUB_DEVICE inline
62  T Sum(T input)
63  {
64  base_type::reduce(input, input, temp_storage_);
65  return input;
66  }
67 
68  HIPCUB_DEVICE inline
69  T Sum(T input, int valid_items)
70  {
71  base_type::reduce(input, input, valid_items, temp_storage_);
72  return input;
73  }
74 
75  template<typename FlagT>
76  HIPCUB_DEVICE inline
77  T HeadSegmentedSum(T input, FlagT head_flag)
78  {
79  base_type::head_segmented_reduce(input, input, head_flag, temp_storage_);
80  return input;
81  }
82 
83  template<typename FlagT>
84  HIPCUB_DEVICE inline
85  T TailSegmentedSum(T input, FlagT tail_flag)
86  {
87  base_type::tail_segmented_reduce(input, input, tail_flag, temp_storage_);
88  return input;
89  }
90 
91  template<typename ReduceOp>
92  HIPCUB_DEVICE inline
93  T Reduce(T input, ReduceOp reduce_op)
94  {
95  base_type::reduce(input, input, temp_storage_, reduce_op);
96  return input;
97  }
98 
99  template<typename ReduceOp>
100  HIPCUB_DEVICE inline
101  T Reduce(T input, ReduceOp reduce_op, int valid_items)
102  {
103  base_type::reduce(input, input, valid_items, temp_storage_, reduce_op);
104  return input;
105  }
106 
107  template<typename ReduceOp, typename FlagT>
108  HIPCUB_DEVICE inline
109  T HeadSegmentedReduce(T input, FlagT head_flag, ReduceOp reduce_op)
110  {
111  base_type::head_segmented_reduce(
112  input, input, head_flag, temp_storage_, reduce_op
113  );
114  return input;
115  }
116 
117  template<typename ReduceOp, typename FlagT>
118  HIPCUB_DEVICE inline
119  T TailSegmentedReduce(T input, FlagT tail_flag, ReduceOp reduce_op)
120  {
121  base_type::tail_segmented_reduce(
122  input, input, tail_flag, temp_storage_, reduce_op
123  );
124  return input;
125  }
126 };
127 
128 END_HIPCUB_NAMESPACE
129 
130 #endif // HIPCUB_ROCPRIM_WARP_WARP_REDUCE_HPP_
hipcub::WarpReduce
Definition: warp_reduce.hpp:47
hipcub::Sum
Definition: thread_operators.hpp:76