/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-hipcub/checkouts/docs-5.0.2/hipcub/include/hipcub/backend/rocprim/device/device_scan.hpp Source File

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

device_scan.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_DEVICE_DEVICE_SCAN_HPP_

31 #define HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_

32 

33 #include <iostream>

34 #include "../../../config.hpp"

35 

36 #include "../thread/thread_operators.hpp"

37 

38 #include <rocprim/device/device_scan.hpp>

39 BEGIN_HIPCUB_NAMESPACE

40 

41 class DeviceScan

42 {

43 public:

44  template <

45  typename InputIteratorT,

46  typename OutputIteratorT

47  >

48  HIPCUB_RUNTIME_FUNCTION static

49  hipError_t InclusiveSum(void *d_temp_storage,

50  size_t &temp_storage_bytes,

51  InputIteratorT d_in,

52  OutputIteratorT d_out,

53  size_t num_items,

54  hipStream_t stream = 0,

55  bool debug_synchronous = false)

56  {

57  return InclusiveScan(

58  d_temp_storage, temp_storage_bytes,

59  d_in, d_out, ::hipcub::Sum(), num_items,

60  stream, debug_synchronous

61  );

62  }

63 

64  template <

65  typename InputIteratorT,

66  typename OutputIteratorT,

67  typename ScanOpT

68  >

69  HIPCUB_RUNTIME_FUNCTION static

70  hipError_t InclusiveScan(void *d_temp_storage,

71  size_t &temp_storage_bytes,

72  InputIteratorT d_in,

73  OutputIteratorT d_out,

74  ScanOpT scan_op,

75  size_t num_items,

76  hipStream_t stream = 0,

77  bool debug_synchronous = false)

78  {

79  return ::rocprim::inclusive_scan(

80  d_temp_storage, temp_storage_bytes,

81  d_in, d_out, num_items,

82  ::hipcub::detail::convert_result_type<InputIteratorT, OutputIteratorT>(scan_op),

83  stream, debug_synchronous

84  );

85  }

86 

87  template <

88  typename InputIteratorT,

89  typename OutputIteratorT

90  >

91  HIPCUB_RUNTIME_FUNCTION static

92  hipError_t ExclusiveSum(void *d_temp_storage,

93  size_t &temp_storage_bytes,

94  InputIteratorT d_in,

95  OutputIteratorT d_out,

96  size_t num_items,

97  hipStream_t stream = 0,

98  bool debug_synchronous = false)

99  {

100  using T = typename std::iterator_traits<InputIteratorT>::value_type;

101  return ExclusiveScan(

102  d_temp_storage, temp_storage_bytes,

103  d_in, d_out, ::hipcub::Sum(), T(0), num_items,

104  stream, debug_synchronous

105  );

106  }

107 

108  template <

109  typename InputIteratorT,

110  typename OutputIteratorT,

111  typename ScanOpT,

112  typename InitValueT

113  >

114  HIPCUB_RUNTIME_FUNCTION static

115  hipError_t ExclusiveScan(void *d_temp_storage,

116  size_t &temp_storage_bytes,

117  InputIteratorT d_in,

118  OutputIteratorT d_out,

119  ScanOpT scan_op,

120  InitValueT init_value,

121  size_t num_items,

122  hipStream_t stream = 0,

123  bool debug_synchronous = false)

124  {

125  return ::rocprim::exclusive_scan(

126  d_temp_storage, temp_storage_bytes,

127  d_in, d_out, init_value, num_items,

128  ::hipcub::detail::convert_result_type<InputIteratorT, OutputIteratorT>(scan_op),

129  stream, debug_synchronous

130  );

131  }

132 };

133 

134 END_HIPCUB_NAMESPACE

135 

136 #endif // HIPCUB_ROCPRIM_DEVICE_DEVICE_SCAN_HPP_

hipcub::Sum

Definition: thread_operators.hpp:76