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

• develop
• include
• ck_tile
• ops
• gemm
• pipeline

tile_gemm_traits.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

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

3 

4 #pragma once

5 

6 #include "ck_tile/core.hpp"

7 

8 namespace ck_tile {

9 

10 template <bool kPadM_,

11  bool kPadN_,

12  bool kPadK_,

13  typename AsLayout_,

14  typename BsLayout_,

15  typename CLayout_,

16  index_t NumWaveGroups_ = 1>

17 struct TileGemmTraits

18 {

19  static constexpr bool kPadM = kPadM_;

20  static constexpr bool kPadN = kPadN_;

21  static constexpr bool kPadK = kPadK_;

22 

23  // TODO this can't be hardcoded here! Should be in policy!

24  static constexpr int _VectorSize = 16;

25 

26  using AsLayout = AsLayout_;

27  using BsLayout = BsLayout_;

28  using CLayout = CLayout_;

29 

30  static constexpr bool TransposeC = false;

31  static constexpr bool UseStructuredSparsity = false;

32  static constexpr index_t NumWaveGroups = NumWaveGroups_;

33 };

34 

35 template <bool kPadM_,

36  bool kPadN_,

37  bool kPadK_,

38  bool DoubleSmemBuffer_,

39  typename AsLayout_,

40  typename BsLayout_,

41  typename CLayout_,

42  bool TransposeC_ = false,

43  bool UseStructuredSparsity_ = false,

44  bool UsePersistentKernel_ = false,

45  index_t NumWaveGroups_ = 1,

46  bool Preshuffle_ = false>

47 struct TileGemmUniversalTraits

48 {

49  static constexpr bool kPadM = kPadM_;

50  static constexpr bool kPadN = kPadN_;

51  static constexpr bool kPadK = kPadK_;

52  static constexpr int _VectorSize = 16;

53  static constexpr bool DoubleSmemBuffer = DoubleSmemBuffer_;

54 

55  using AsLayout = AsLayout_;

56  using BsLayout = BsLayout_;

57  using CLayout = CLayout_;

58 

59  static constexpr bool TransposeC = TransposeC_;

60  static constexpr bool UseStructuredSparsity = UseStructuredSparsity_;

61  static constexpr bool UsePersistentKernel = UsePersistentKernel_;

62  static constexpr index_t NumWaveGroups = NumWaveGroups_;

63  static constexpr bool Preshuffle = Preshuffle_;

64 };

65 

66 template <bool kPadM_,

67  bool kPadN_,

68  bool kPadK_,

69  bool DoubleSmemBuffer_,

70  typename AsLayout_,

71  typename BsLayout_,

72  typename CLayout_,

73  bool TransposeC_ = false,

74  bool UseStructuredSparsity_ = false>

75 using PersistentTileGemmUniversalTraits = TileGemmUniversalTraits<kPadM_,

76  kPadN_,

77  kPadK_,

78  DoubleSmemBuffer_,

79  AsLayout_,

80  BsLayout_,

81  CLayout_,

82  TransposeC_,

83  UseStructuredSparsity_,

84  true>;

85 

86 } // namespace ck_tile

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::TileGemmTraits

Definition: tile_gemm_traits.hpp:18

ck_tile::TileGemmTraits::NumWaveGroups

static constexpr index_t NumWaveGroups

Definition: tile_gemm_traits.hpp:32

ck_tile::TileGemmTraits::CLayout

CLayout_ CLayout

Definition: tile_gemm_traits.hpp:28

ck_tile::TileGemmTraits::TransposeC

static constexpr bool TransposeC

Definition: tile_gemm_traits.hpp:30

ck_tile::TileGemmTraits::AsLayout

AsLayout_ AsLayout

Definition: tile_gemm_traits.hpp:26

ck_tile::TileGemmTraits::kPadM

static constexpr bool kPadM

Definition: tile_gemm_traits.hpp:19

ck_tile::TileGemmTraits::_VectorSize

static constexpr int _VectorSize

Definition: tile_gemm_traits.hpp:24

ck_tile::TileGemmTraits::kPadN

static constexpr bool kPadN

Definition: tile_gemm_traits.hpp:20

ck_tile::TileGemmTraits::UseStructuredSparsity

static constexpr bool UseStructuredSparsity

Definition: tile_gemm_traits.hpp:31

ck_tile::TileGemmTraits::kPadK

static constexpr bool kPadK

Definition: tile_gemm_traits.hpp:21

ck_tile::TileGemmTraits::BsLayout

BsLayout_ BsLayout

Definition: tile_gemm_traits.hpp:27

ck_tile::TileGemmUniversalTraits

Definition: tile_gemm_traits.hpp:48

ck_tile::TileGemmUniversalTraits::kPadN

static constexpr bool kPadN

Definition: tile_gemm_traits.hpp:50

ck_tile::TileGemmUniversalTraits::TransposeC

static constexpr bool TransposeC

Definition: tile_gemm_traits.hpp:59

ck_tile::TileGemmUniversalTraits::UseStructuredSparsity

static constexpr bool UseStructuredSparsity

Definition: tile_gemm_traits.hpp:60

ck_tile::TileGemmUniversalTraits::DoubleSmemBuffer

static constexpr bool DoubleSmemBuffer

Definition: tile_gemm_traits.hpp:53

ck_tile::TileGemmUniversalTraits::kPadM

static constexpr bool kPadM

Definition: tile_gemm_traits.hpp:49

ck_tile::TileGemmUniversalTraits::_VectorSize

static constexpr int _VectorSize

Definition: tile_gemm_traits.hpp:52

ck_tile::TileGemmUniversalTraits::CLayout

CLayout_ CLayout

Definition: tile_gemm_traits.hpp:57

ck_tile::TileGemmUniversalTraits::NumWaveGroups

static constexpr index_t NumWaveGroups

Definition: tile_gemm_traits.hpp:62

ck_tile::TileGemmUniversalTraits::kPadK

static constexpr bool kPadK

Definition: tile_gemm_traits.hpp:51

ck_tile::TileGemmUniversalTraits::BsLayout

BsLayout_ BsLayout

Definition: tile_gemm_traits.hpp:56

ck_tile::TileGemmUniversalTraits::UsePersistentKernel

static constexpr bool UsePersistentKernel

Definition: tile_gemm_traits.hpp:61

ck_tile::TileGemmUniversalTraits::Preshuffle

static constexpr bool Preshuffle

Definition: tile_gemm_traits.hpp:63

ck_tile::TileGemmUniversalTraits::AsLayout

AsLayout_ AsLayout

Definition: tile_gemm_traits.hpp:55