/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.3.3/hipcub/include/hipcub/backend/rocprim/thread/thread_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
• thread

thread_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_THREAD_THREAD_REDUCE_HPP_

31 #define HIPCUB_ROCPRIM_THREAD_THREAD_REDUCE_HPP_

32 

33 BEGIN_HIPCUB_NAMESPACE

34 

36 namespace internal {

37 

38 template <

39  int LENGTH,

40  typename T,

41  typename ReductionOp,

42  bool NoPrefix = false>

43 __device__ __forceinline__ T ThreadReduce(

44  T* input,

45  ReductionOp reduction_op,

46  T prefix = T(0))

47 {

48  T retval;

49  if(NoPrefix)

50  retval = input[0];

51  else

52  retval = prefix;

53 

54  #pragma unroll

55  for (int i = 0 + NoPrefix; i < LENGTH; ++i)

56  retval = reduction_op(retval, input[i]);

57 

58  return retval;

59 }

60 

61 template <

62  int LENGTH,

63  typename T,

64  typename ReductionOp>

65 __device__ __forceinline__ T ThreadReduce(

66  T (&input)[LENGTH],

67  ReductionOp reduction_op,

68  T prefix)

69 {

70  return ThreadReduce<LENGTH, false>((T*)input, reduction_op, prefix);

71 }

72 

73 template <

74  int LENGTH,

75  typename T,

76  typename ReductionOp>

77 __device__ __forceinline__ T ThreadReduce(

78  T (&input)[LENGTH],

79  ReductionOp reduction_op)

80 {

81  return ThreadReduce<LENGTH, true>((T*)input, reduction_op);

82 }

83 

84 }

85 

86 END_HIPCUB_NAMESPACE

87 

88 #endif

internal

Internal namespace (to prevent ADL mishaps between static functions when mixing different CUB install...

Definition: thread_reduce.hpp:36