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

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

block_store_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_STORE_FUNC_HPP_

31 #define HIPCUB_ROCPRIM_BLOCK_BLOCK_STORE_FUNC_HPP_

32 

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

34 

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

36 

37 BEGIN_HIPCUB_NAMESPACE

38 

39 template<

40  typename T,

41  int ITEMS_PER_THREAD,

42  typename OutputIteratorT

43 >

44 HIPCUB_DEVICE inline

45 void StoreDirectBlocked(int linear_id,

46  OutputIteratorT block_iter,

47  T (&items)[ITEMS_PER_THREAD])

48 {

49  ::rocprim::block_store_direct_blocked(

50  linear_id, block_iter, items

51  );

52 }

53 

54 template<

55  typename T,

56  int ITEMS_PER_THREAD,

57  typename OutputIteratorT

58 >

59 HIPCUB_DEVICE inline

60 void StoreDirectBlocked(int linear_id,

61  OutputIteratorT block_iter,

62  T (&items)[ITEMS_PER_THREAD],

63  int valid_items)

64 {

65  ::rocprim::block_store_direct_blocked(

66  linear_id, block_iter, items, valid_items

67  );

68 }

69 

70 template <

71  typename T,

72  int ITEMS_PER_THREAD

73 >

74 HIPCUB_DEVICE inline

75 void StoreDirectBlockedVectorized(int linear_id,

76  T* block_iter,

77  T (&items)[ITEMS_PER_THREAD])

78 {

79  ::rocprim::block_store_direct_blocked_vectorized(

80  linear_id, block_iter, items

81  );

82 }

83 

84 template<

85  int BLOCK_THREADS,

86  typename T,

87  int ITEMS_PER_THREAD,

88  typename OutputIteratorT

89 >

90 HIPCUB_DEVICE inline

91 void StoreDirectStriped(int linear_id,

92  OutputIteratorT block_iter,

93  T (&items)[ITEMS_PER_THREAD])

94 {

95  ::rocprim::block_store_direct_striped<BLOCK_THREADS>(

96  linear_id, block_iter, items

97  );

98 }

99 

100 template<

101  int BLOCK_THREADS,

102  typename T,

103  int ITEMS_PER_THREAD,

104  typename OutputIteratorT

105 >

106 HIPCUB_DEVICE inline

107 void StoreDirectStriped(int linear_id,

108  OutputIteratorT block_iter,

109  T (&items)[ITEMS_PER_THREAD],

110  int valid_items)

111 {

112  ::rocprim::block_store_direct_striped<BLOCK_THREADS>(

113  linear_id, block_iter, items, valid_items

114  );

115 }

116 

117 template<

118  typename T,

119  int ITEMS_PER_THREAD,

120  typename OutputIteratorT

121 >

122 HIPCUB_DEVICE inline

123 void StoreDirectWarpStriped(int linear_id,

124  OutputIteratorT block_iter,

125  T (&items)[ITEMS_PER_THREAD])

126 {

127  ::rocprim::block_store_direct_warp_striped(

128  linear_id, block_iter, items

129  );

130 }

131 

132 template<

133  typename T,

134  int ITEMS_PER_THREAD,

135  typename OutputIteratorT

136 >

137 HIPCUB_DEVICE inline

138 void StoreDirectWarpStriped(int linear_id,

139  OutputIteratorT block_iter,

140  T (&items)[ITEMS_PER_THREAD],

141  int valid_items)

142 {

143  ::rocprim::block_store_direct_warp_striped(

144  linear_id, block_iter, items, valid_items

145  );

146 }

147 

148 END_HIPCUB_NAMESPACE

149 

150 #endif // HIPCUB_ROCPRIM_BLOCK_BLOCK_STORE_FUNC_HPP_