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

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

warp_exchange.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-2021, 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_EXCHANGE_HPP_

31 #define HIPCUB_ROCPRIM_WARP_WARP_EXCHANGE_HPP_

32 

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

34 #include "../util_type.hpp"

35 

36 #include <rocprim/warp/warp_exchange.hpp>

37 

38 BEGIN_HIPCUB_NAMESPACE

39 

40 template <

41  typename InputT,

42  int ITEMS_PER_THREAD,

43  int LOGICAL_WARP_THREADS = HIPCUB_DEVICE_WARP_THREADS,

44  int ARCH = HIPCUB_ARCH

45 >

46 class WarpExchange

47 {

48  using base_type = typename rocprim::warp_exchange<InputT, ITEMS_PER_THREAD, LOGICAL_WARP_THREADS>;

49 

50 public:

51  using TempStorage = typename base_type::storage_type;

52 

53 private:

54  TempStorage &temp_storage;

55 

56 public:

57  WarpExchange() = delete;

58 

59  explicit HIPCUB_DEVICE __forceinline__

60  WarpExchange(TempStorage &temp_storage) :

61  temp_storage(temp_storage)

62  {

63  }

64 

65  template <typename OutputT>

66  HIPCUB_DEVICE __forceinline__

67  void BlockedToStriped(

68  const InputT (&input_items)[ITEMS_PER_THREAD],

69  OutputT (&output_items)[ITEMS_PER_THREAD])

70  {

71  base_type rocprim_warp_exchange;

72  rocprim_warp_exchange.blocked_to_striped(input_items, output_items, temp_storage);

73  }

74 

75  template <typename OutputT>

76  HIPCUB_DEVICE __forceinline__

77  void StripedToBlocked(

78  const InputT (&input_items)[ITEMS_PER_THREAD],

79  OutputT (&output_items)[ITEMS_PER_THREAD])

80  {

81  base_type rocprim_warp_exchange;

82  rocprim_warp_exchange.striped_to_blocked(input_items, output_items, temp_storage);

83  }

84 

85  template <typename OffsetT>

86  HIPCUB_DEVICE __forceinline__

87  void ScatterToStriped(

88  InputT (&items)[ITEMS_PER_THREAD],

89  OffsetT (&ranks)[ITEMS_PER_THREAD])

90  {

91  ScatterToStriped(items, items, ranks);

92  }

93 

94  template <typename OutputT,

95  typename OffsetT>

96  HIPCUB_DEVICE __forceinline__

97  void ScatterToStriped(

98  const InputT (&input_items)[ITEMS_PER_THREAD],

99  OutputT (&output_items)[ITEMS_PER_THREAD],

100  OffsetT (&ranks)[ITEMS_PER_THREAD])

101  {

102  base_type rocprim_warp_exchange;

103  rocprim_warp_exchange.scatter_to_striped(input_items, output_items, ranks, temp_storage);

104  }

105 };

106 

107 END_HIPCUB_NAMESPACE

108 

109 #endif // HIPCUB_ROCPRIM_WARP_WARP_EXCHANGE_HPP_

hipcub::WarpExchange

Definition: warp_exchange.hpp:47