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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/docs-7.0.0/include/ck_tile/ops/gemm/pipeline/tile_gemm_traits.hpp Source File
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 ALayout_,
14  typename BLayout_,
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 ALayout = ALayout_;
27  using BLayout = BLayout_;
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 ALayout_,
40  typename BLayout_,
41  typename CLayout_,
42  bool TransposeC_ = false,
43  bool UseStructuredSparsity_ = false,
44  bool UsePersistentKernel_ = false,
45  index_t NumWaveGroups_ = 1>
46 struct TileGemmUniversalTraits
47 {
48  static constexpr bool kPadM = kPadM_;
49  static constexpr bool kPadN = kPadN_;
50  static constexpr bool kPadK = kPadK_;
51 
52  static constexpr bool DoubleSmemBuffer = DoubleSmemBuffer_;
53 
54  using ALayout = ALayout_;
55  using BLayout = BLayout_;
56  using CLayout = CLayout_;
57 
58  static constexpr bool TransposeC = TransposeC_;
59  static constexpr bool UseStructuredSparsity = UseStructuredSparsity_;
60  static constexpr bool UsePersistentKernel = UsePersistentKernel_;
61  static constexpr index_t NumWaveGroups = NumWaveGroups_;
62 };
63 
64 template <bool kPadM_,
65  bool kPadN_,
66  bool kPadK_,
67  bool DoubleSmemBuffer_,
68  typename ALayout_,
69  typename BLayout_,
70  typename CLayout_,
71  bool TransposeC_ = false,
72  bool UseStructuredSparsity_ = false>
73 using PersistentTileGemmUniversalTraits = TileGemmUniversalTraits<kPadM_,
74  kPadN_,
75  kPadK_,
76  DoubleSmemBuffer_,
77  ALayout_,
78  BLayout_,
79  CLayout_,
80  TransposeC_,
81  UseStructuredSparsity_,
82  true>;
83 
84 } // namespace ck_tile
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::UseStructuredSparsity
static constexpr bool UseStructuredSparsity
Definition: tile_gemm_traits.hpp:31
ck_tile::TileGemmTraits::ALayout
ALayout_ ALayout
Definition: tile_gemm_traits.hpp:26
ck_tile::TileGemmTraits::kPadM
static constexpr bool kPadM
Definition: tile_gemm_traits.hpp:19
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::_VectorSize
static constexpr int _VectorSize
Definition: tile_gemm_traits.hpp:24
ck_tile::TileGemmTraits::kPadK
static constexpr bool kPadK
Definition: tile_gemm_traits.hpp:21
ck_tile::TileGemmTraits::kPadN
static constexpr bool kPadN
Definition: tile_gemm_traits.hpp:20
ck_tile::TileGemmTraits::BLayout
BLayout_ BLayout
Definition: tile_gemm_traits.hpp:27
ck_tile::TileGemmUniversalTraits
Definition: tile_gemm_traits.hpp:47
ck_tile::TileGemmUniversalTraits::TransposeC
static constexpr bool TransposeC
Definition: tile_gemm_traits.hpp:58
ck_tile::TileGemmUniversalTraits::DoubleSmemBuffer
static constexpr bool DoubleSmemBuffer
Definition: tile_gemm_traits.hpp:52
ck_tile::TileGemmUniversalTraits::UsePersistentKernel
static constexpr bool UsePersistentKernel
Definition: tile_gemm_traits.hpp:60
ck_tile::TileGemmUniversalTraits::UseStructuredSparsity
static constexpr bool UseStructuredSparsity
Definition: tile_gemm_traits.hpp:59
ck_tile::TileGemmUniversalTraits::ALayout
ALayout_ ALayout
Definition: tile_gemm_traits.hpp:54
ck_tile::TileGemmUniversalTraits::BLayout
BLayout_ BLayout
Definition: tile_gemm_traits.hpp:55
ck_tile::TileGemmUniversalTraits::kPadK
static constexpr bool kPadK
Definition: tile_gemm_traits.hpp:50
ck_tile::TileGemmUniversalTraits::kPadN
static constexpr bool kPadN
Definition: tile_gemm_traits.hpp:49
ck_tile::TileGemmUniversalTraits::kPadM
static constexpr bool kPadM
Definition: tile_gemm_traits.hpp:48
ck_tile::TileGemmUniversalTraits::CLayout
CLayout_ CLayout
Definition: tile_gemm_traits.hpp:56
ck_tile::TileGemmUniversalTraits::NumWaveGroups
static constexpr index_t NumWaveGroups
Definition: tile_gemm_traits.hpp:61