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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm_quant/pipeline/tile_gemm_quant_traits.hpp Source File
tile_gemm_quant_traits.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include "ck_tile/core.hpp"
7 #include <cstdint>
8 
9 namespace ck_tile {
10 
11 enum struct QuantType : std::uint16_t
12 {
13  AQuantGrouped = 0,
14  BQuantGrouped = 1,
15  RowColQuant = 2,
16  TensorQuant = 3
17 };
18 
19 inline std::string quant_type_to_string(QuantType quant_type)
20 {
21  switch(quant_type)
22  {
23  case QuantType::AQuantGrouped: return "AQuantGrouped";
24  case QuantType::BQuantGrouped: return "BQuantGrouped";
25  case QuantType::RowColQuant: return "RowColQuant";
26  case QuantType::TensorQuant: return "TensorQuant";
27  default: return "Unknown";
28  }
29 }
30 
31 template <bool kPadM_,
32  bool kPadN_,
33  bool kPadK_,
34  bool PreshuffleQuant_,
35  bool PreshuffleB_,
36  typename ALayout_,
37  typename BLayout_,
38  typename CLayout_,
39  QuantType QuantType_,
40  typename AQLayout_ = ALayout_,
41  typename BQLayout_ = BLayout_,
42  bool DoubleSmemBuffer_ = false,
43  bool UsePersistentKernel_ = false>
45 {
46  static constexpr bool kPadM = kPadM_;
47  static constexpr bool kPadN = kPadN_;
48  static constexpr bool kPadK = kPadK_;
49 
50  static constexpr QuantType kQuantType = QuantType_;
51 
52  static constexpr int _VectorSize = 16;
53  static constexpr bool DoubleSmemBuffer = DoubleSmemBuffer_;
54 
55  using ALayout = ALayout_;
56  using BLayout = BLayout_;
57  using CLayout = CLayout_;
58  using AQLayout = AQLayout_;
59  using BQLayout = BQLayout_;
60 
61  // TODO: It should be replaced to single value
62  using AsLayout = ALayout_;
63  using BsLayout = BLayout_;
64 
65  static constexpr bool TransposeC = false;
66  static constexpr bool UseStructuredSparsity = false;
67  static constexpr index_t NumWaveGroups = 1;
68  static constexpr bool UsePersistentKernel = UsePersistentKernel_;
69 
70  static constexpr bool PreshuffleQuant = PreshuffleQuant_;
71  static constexpr bool PreshuffleB = PreshuffleB_;
72 };
73 
74 } // namespace ck_tile
Definition: cluster_descriptor.hpp:13
std::string quant_type_to_string(QuantType quant_type)
Definition: tile_gemm_quant_traits.hpp:19
int32_t index_t
Definition: integer.hpp:9
QuantType
Definition: tile_gemm_quant_traits.hpp:12
unsigned short uint16_t
Definition: stdint.h:125
Definition: tile_gemm_quant_traits.hpp:45
static constexpr int _VectorSize
Definition: tile_gemm_quant_traits.hpp:52
static constexpr index_t NumWaveGroups
Definition: tile_gemm_quant_traits.hpp:67
static constexpr bool UsePersistentKernel
Definition: tile_gemm_quant_traits.hpp:68
static constexpr bool UseStructuredSparsity
Definition: tile_gemm_quant_traits.hpp:66
ALayout_ AsLayout
Definition: tile_gemm_quant_traits.hpp:62
static constexpr bool PreshuffleQuant
Definition: tile_gemm_quant_traits.hpp:70
BLayout_ BLayout
Definition: tile_gemm_quant_traits.hpp:56
static constexpr bool kPadN
Definition: tile_gemm_quant_traits.hpp:47
static constexpr bool kPadM
Definition: tile_gemm_quant_traits.hpp:46
static constexpr bool PreshuffleB
Definition: tile_gemm_quant_traits.hpp:71
AQLayout_ AQLayout
Definition: tile_gemm_quant_traits.hpp:58
static constexpr QuantType kQuantType
Definition: tile_gemm_quant_traits.hpp:50
CLayout_ CLayout
Definition: tile_gemm_quant_traits.hpp:57
static constexpr bool TransposeC
Definition: tile_gemm_quant_traits.hpp:65
BQLayout_ BQLayout
Definition: tile_gemm_quant_traits.hpp:59
BLayout_ BsLayout
Definition: tile_gemm_quant_traits.hpp:63
ALayout_ ALayout
Definition: tile_gemm_quant_traits.hpp:55
static constexpr bool DoubleSmemBuffer
Definition: tile_gemm_quant_traits.hpp:53
static constexpr bool kPadK
Definition: tile_gemm_quant_traits.hpp:48