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

• develop
• include
• ck_tile
• ops
• flatmm
• pipeline

tile_flatmm_shape.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/host/concat.hpp"

8 

9 namespace ck_tile {

10 

11 template <typename BlockTile_, typename BlockWarps_, typename WarpTile_>

12 struct TileFlatmmShape

13 {

14  using BlockTile = remove_cvref_t<BlockTile_>;

15  using BlockWarps = remove_cvref_t<BlockWarps_>;

16  using WarpTile = remove_cvref_t<WarpTile_>;

17 

18  static constexpr auto idxM = number<0>{};

19  static constexpr auto idxN = number<1>{};

20  static constexpr auto idxK = number<2>{};

21 

22  static constexpr index_t NumWarps = reduce_on_sequence(BlockWarps{}, multiplies{}, number<1>{});

23 

24  static constexpr index_t kM = BlockTile::at(idxM);

25  static constexpr index_t kN = BlockTile::at(idxN);

26  static constexpr index_t kK = BlockTile::at(idxK);

27 

28  static constexpr index_t flatNPerWarp = BlockWarps::at(idxN);

29  static constexpr index_t flatKPerWarp = WarpTile::at(idxK) * WarpTile::at(idxN);

30  static constexpr index_t flatKPerBlock = flatKPerWarp * kK / WarpTile::at(idxK);

31 

32  static constexpr bool PermuteA = false;

33  static constexpr bool PermuteB = false;

34 

35  CK_TILE_HOST static std::string GetName()

36  {

37  // clang-format off

38  return concat('_', "tile_flatmm_shape",

39  concat('x', kM, kN, kK, NumWarps),

40  concat('x', BlockWarps::at(idxM), BlockWarps::at(idxN), BlockWarps::at(idxK)),

41  concat('x', (WarpTile::at(idxM)), WarpTile::at(idxN), WarpTile::at(idxK)));

42  // clang-format on

43  }

44 };

45 

46 } // namespace ck_tile

concat.hpp

CK_TILE_HOST

#define CK_TILE_HOST

Definition: config.hpp:40

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::concat

auto concat(const Ts &... xs) -> std::enable_if_t<!AllConvertibleToStringView< Ts... >, std::string >

Definition: concat.hpp:43

ck_tile::remove_cvref_t

remove_cv_t< std::remove_reference_t< T > > remove_cvref_t

Definition: type_traits.hpp:21

ck_tile::reduce_on_sequence

constexpr CK_TILE_HOST_DEVICE index_t reduce_on_sequence(Seq, Reduce f, number< Init >)

Definition: sequence.hpp:979

ck_tile::TileFlatmmShape

Definition: tile_flatmm_shape.hpp:13

ck_tile::TileFlatmmShape::flatKPerWarp

static constexpr index_t flatKPerWarp

Definition: tile_flatmm_shape.hpp:29

ck_tile::TileFlatmmShape::kK

static constexpr index_t kK

Definition: tile_flatmm_shape.hpp:26

ck_tile::TileFlatmmShape::flatKPerBlock

static constexpr index_t flatKPerBlock

Definition: tile_flatmm_shape.hpp:30

ck_tile::TileFlatmmShape::idxK

static constexpr auto idxK

Definition: tile_flatmm_shape.hpp:20

ck_tile::TileFlatmmShape::NumWarps

static constexpr index_t NumWarps

Definition: tile_flatmm_shape.hpp:22

ck_tile::TileFlatmmShape::kN

static constexpr index_t kN

Definition: tile_flatmm_shape.hpp:25

ck_tile::TileFlatmmShape::PermuteB

static constexpr bool PermuteB

Definition: tile_flatmm_shape.hpp:33

ck_tile::TileFlatmmShape::idxN

static constexpr auto idxN

Definition: tile_flatmm_shape.hpp:19

ck_tile::TileFlatmmShape::PermuteA

static constexpr bool PermuteA

Definition: tile_flatmm_shape.hpp:32

ck_tile::TileFlatmmShape::flatNPerWarp

static constexpr index_t flatNPerWarp

Definition: tile_flatmm_shape.hpp:28

ck_tile::TileFlatmmShape::idxM

static constexpr auto idxM

Definition: tile_flatmm_shape.hpp:18

ck_tile::TileFlatmmShape::WarpTile

remove_cvref_t< WarpTile_ > WarpTile

Definition: tile_flatmm_shape.hpp:16

ck_tile::TileFlatmmShape::BlockTile

remove_cvref_t< BlockTile_ > BlockTile

Definition: tile_flatmm_shape.hpp:14

ck_tile::TileFlatmmShape::BlockWarps

remove_cvref_t< BlockWarps_ > BlockWarps

Definition: tile_flatmm_shape.hpp:15

ck_tile::TileFlatmmShape::GetName

static CK_TILE_HOST std::string GetName()

Definition: tile_flatmm_shape.hpp:35

ck_tile::TileFlatmmShape::kM

static constexpr index_t kM

Definition: tile_flatmm_shape.hpp:24

ck_tile::constant

Definition: integral_constant.hpp:13

ck_tile::multiplies

Definition: math.hpp:98