/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/reduce/pipeline/reduce2d_default_policy.hpp Source File

• develop
• include
• ck_tile
• ops
• reduce
• pipeline

reduce2d_default_policy.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.

3 

4 #pragma once

5 

6 #include "ck_tile/core.hpp"

7 #include "ck_tile/ops/reduce/block/block_reduce2d_problem.hpp"

8 #include "ck_tile/ops/reduce/block/block_reduce2d.hpp"

9 

10 namespace ck_tile {

11 

12 struct Reduce2dDefaultPolicy

13 {

14  template <typename Problem>

15  CK_TILE_DEVICE static constexpr auto MakeXBlockTileDistribution()

16  {

17  using S = typename Problem::BlockShape;

18  return make_static_tile_distribution(

19  tile_distribution_encoding<

20  sequence<>,

21  tuple<

22  sequence<S::Repeat_M, S::WarpPerBlock_M, S::ThreadPerWarp_M, S::ThreadTile_M>,

23  sequence<S::Repeat_N, S::WarpPerBlock_N, S::ThreadPerWarp_N, S::ThreadTile_N>>,

24  tuple<sequence<1, 2>, sequence<1, 2>>,

25  tuple<sequence<1, 1>, sequence<2, 2>>,

26  sequence<1, 1, 2, 2>,

27  sequence<0, 3, 0, 3>>{});

28  }

29 

30  template <typename Problem>

31  CK_TILE_HOST_DEVICE static constexpr auto GetBlockReduce2d()

32  {

33  using P_ = BlockReduce2dProblem<typename Problem::XDataType,

34  typename Problem::ComputeDataType,

35  typename Problem::BlockShape>;

36  return BlockReduce2d<P_>{};

37  }

38 

39  template <typename Problem>

40  CK_TILE_HOST_DEVICE static constexpr auto GetBlockReduce2dSync()

41  {

42  using P_ = BlockReduce2dProblem<typename Problem::XDataType,

43  typename Problem::ComputeDataType,

44  typename Problem::BlockShape>;

45  return BlockReduce2dSync<P_>{};

46  }

47 

48  template <typename Problem>

49  CK_TILE_HOST_DEVICE static constexpr auto GetBlockReduce2dCrossWarpSync()

50  {

51  using P_ = BlockReduce2dProblem<typename Problem::XDataType,

52  typename Problem::ComputeDataType,

53  typename Problem::BlockShape>;

54  return BlockReduce2dCrossWarpSync<P_>{};

55  }

56 

57  template <typename Problem>

58  CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize()

59  {

60  if constexpr(Problem::kNeedCrossWarpSync)

61  {

62  using P_ = BlockReduce2dProblem<typename Problem::XDataType,

63  typename Problem::ComputeDataType,

64  typename Problem::BlockShape>;

65 

66  using block_reduce2d = BlockReduce2d<P_>;

67  using x_block_tile =

68  decltype(make_static_distributed_tensor<typename Problem::XDataType>(

69  MakeXBlockTileDistribution<Problem>()));

70  using y_block_tile = decltype(block_reduce2d::template MakeYBlockTile<x_block_tile>());

71 

72  return GetBlockReduce2dCrossWarpSync<Problem>().template GetSmemSize<y_block_tile>();

73  }

74  else

75  {

76  return 1; // zero size arrays are an extension

77  }

78  }

79 };

80 } // namespace ck_tile

block_reduce2d.hpp

block_reduce2d_problem.hpp

CK_TILE_DEVICE

#define CK_TILE_DEVICE

Definition: config.hpp:41

CK_TILE_HOST_DEVICE

#define CK_TILE_HOST_DEVICE

Definition: config.hpp:42

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::make_static_tile_distribution

constexpr CK_TILE_HOST_DEVICE auto make_static_tile_distribution(StaticTileDistributionEncoding_)

Definition: tile_distribution.hpp:480

ck_tile::BlockReduce2dCrossWarpSync

Definition: block_reduce2d.hpp:200

ck_tile::BlockReduce2d

Definition: block_reduce2d.hpp:45

ck_tile::BlockReduce2dProblem

Definition: block_reduce2d_problem.hpp:12

ck_tile::BlockReduce2dSync

Definition: block_reduce2d.hpp:135

ck_tile::Reduce2dDefaultPolicy

Definition: reduce2d_default_policy.hpp:13

ck_tile::Reduce2dDefaultPolicy::GetBlockReduce2dSync

static constexpr CK_TILE_HOST_DEVICE auto GetBlockReduce2dSync()

Definition: reduce2d_default_policy.hpp:40

ck_tile::Reduce2dDefaultPolicy::GetSmemSize

static constexpr CK_TILE_HOST_DEVICE index_t GetSmemSize()

Definition: reduce2d_default_policy.hpp:58

ck_tile::Reduce2dDefaultPolicy::MakeXBlockTileDistribution

static constexpr CK_TILE_DEVICE auto MakeXBlockTileDistribution()

Definition: reduce2d_default_policy.hpp:15

ck_tile::Reduce2dDefaultPolicy::GetBlockReduce2d

static constexpr CK_TILE_HOST_DEVICE auto GetBlockReduce2d()

Definition: reduce2d_default_policy.hpp:31

ck_tile::Reduce2dDefaultPolicy::GetBlockReduce2dCrossWarpSync

static constexpr CK_TILE_HOST_DEVICE auto GetBlockReduce2dCrossWarpSync()

Definition: reduce2d_default_policy.hpp:49

ck_tile::sequence

Definition: sequence.hpp:49

ck_tile::tile_distribution_encoding

Definition: tile_distribution_encoding.hpp:26

ck_tile::tuple

Definition: tuple.hpp:192