/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 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.

2 // SPDX-License-Identifier: MIT

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  int VectorSize_ = 16>

48 struct TileGemmUniversalTraits

49 {

50  static constexpr bool kPadM = kPadM_;

51  static constexpr bool kPadN = kPadN_;

52  static constexpr bool kPadK = kPadK_;

53  static constexpr int _VectorSize = VectorSize_;

54  static constexpr bool DoubleSmemBuffer = DoubleSmemBuffer_;

55 

56  using AsLayout = AsLayout_;

57  using BsLayout = BsLayout_;

58  using CLayout = CLayout_;

59  static constexpr bool TransposeC = TransposeC_;

60 

61  static constexpr bool UseStructuredSparsity = UseStructuredSparsity_;

62  static constexpr bool UsePersistentKernel = UsePersistentKernel_;

63  static constexpr index_t NumWaveGroups = NumWaveGroups_;

64  static constexpr bool Preshuffle = Preshuffle_;

65 };

66 

67 template <bool kPadM_,

68  bool kPadN_,

69  bool kPadK_,

70  bool DoubleSmemBuffer_,

71  typename AsLayout_,

72  typename BsLayout_,

73  typename CLayout_,

74  bool TransposeC_ = false,

75  bool UseStructuredSparsity_ = false>

76 using PersistentTileGemmUniversalTraits = TileGemmUniversalTraits<kPadM_,

77  kPadN_,

78  kPadK_,

79  DoubleSmemBuffer_,

80  AsLayout_,

81  BsLayout_,

82  CLayout_,

83  TransposeC_,

84  UseStructuredSparsity_,

85  true>;

86 

87 } // 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:49

ck_tile::TileGemmUniversalTraits::Preshuffle

static constexpr bool Preshuffle

Definition: tile_gemm_traits.hpp:64

ck_tile::TileGemmUniversalTraits::BsLayout

BsLayout_ BsLayout

Definition: tile_gemm_traits.hpp:57

ck_tile::TileGemmUniversalTraits::NumWaveGroups

static constexpr index_t NumWaveGroups

Definition: tile_gemm_traits.hpp:63

ck_tile::TileGemmUniversalTraits::UsePersistentKernel

static constexpr bool UsePersistentKernel

Definition: tile_gemm_traits.hpp:62

ck_tile::TileGemmUniversalTraits::UseStructuredSparsity

static constexpr bool UseStructuredSparsity

Definition: tile_gemm_traits.hpp:61

ck_tile::TileGemmUniversalTraits::AsLayout

AsLayout_ AsLayout

Definition: tile_gemm_traits.hpp:56

ck_tile::TileGemmUniversalTraits::_VectorSize

static constexpr int _VectorSize

Definition: tile_gemm_traits.hpp:53

ck_tile::TileGemmUniversalTraits::DoubleSmemBuffer

static constexpr bool DoubleSmemBuffer

Definition: tile_gemm_traits.hpp:54

ck_tile::TileGemmUniversalTraits::CLayout

CLayout_ CLayout

Definition: tile_gemm_traits.hpp:58

ck_tile::TileGemmUniversalTraits::kPadK

static constexpr bool kPadK

Definition: tile_gemm_traits.hpp:52

ck_tile::TileGemmUniversalTraits::kPadN

static constexpr bool kPadN

Definition: tile_gemm_traits.hpp:51

ck_tile::TileGemmUniversalTraits::TransposeC

static constexpr bool TransposeC

Definition: tile_gemm_traits.hpp:59

ck_tile::TileGemmUniversalTraits::kPadM

static constexpr bool kPadM

Definition: tile_gemm_traits.hpp:50