"src/vscode:/vscode.git/clone" did not exist on "3f0697ad7061f1d54c3d39b43b229283676f6c54"
gemm_sp.h 3.19 KB
Newer Older
1
2
3
4
5
6
7
8
9
/*!
 * \file tl/op/gemm_sp.h
 * \brief Define gemm_sp operator.
 *
 */

#ifndef TVM_TL_OP_GEMM_SP_H_
#define TVM_TL_OP_GEMM_SP_H_

10
#include "operator.h"
11
12

namespace tvm {
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
/**
 * Lower the GemmSP operator into a TIR statement for the given lowering
 * context.
 *
 * Produces the TIR Stmt that implements this operator using the provided
 * lowering arguments. The `analyzer` is used for arithmetic simplifications and
 * may be null.
 *
 * @param T Lowering context and arguments.
 * @returns A TIR `Stmt` implementing the lowered operator.
 */
/**
 * Infer memory/layout mapping for operands and outputs of this operator.
 *
 * Computes a LayoutMap describing how logical tensor layouts map to physical
 * buffer layouts for the given inference `level`.
 *
 * @param T Layout inference inputs (shapes, buffer info, etc.).
 * @param level Inference granularity/level.
 * @returns A LayoutMap describing inferred layouts.
 */
/**
 * Compute a warp-level partitioning (rows, cols) for the given number of warps.
 *
 * Returns a pair (warps_per_row, warps_per_col) describing how to tile the GEMM
 * across warps for the specified `target`. The optional `maybe_hopper_wgmma`
 * enables target-specific adjustments (e.g., CDNA WG/MMA variants) when set.
 *
 * @param num_warps Total number of warps available for the tile.
 * @param target Target device/architecture used to guide partitioning choices.
 * @param maybe_hopper_wgmma Enable target-specific WG/MMA adjustments when
 * true.
 * @returns Pair<int,int> of (warps_per_row, warps_per_col).
 */
/**
 * Create a copy of this TileOperator node as a TileOperator reference.
 *
 * The returned TileOperator refers to a new node that is a copy of this node.
 *
 * @returns A TileOperator that is a clone of this node.
 */
/**
 * Construct a GemmSP TileOperator from call arguments and a buffer map.
 *
 * @param args Array of PrimExpr specifying call-site arguments for the
 * operator.
 * @param vmap Mapping from buffer names to tir::Buffer objects for
 * operands/outputs.
 */
/**
 * Return the singleton Op descriptor for the GemmSP operator.
 *
 * @returns Reference to the operator's Op registration object.
 */
67
68
69
70
namespace tl {

using namespace tir;

71
class GemmSPNode : public TileOperatorNode {
72
public:
73
74
  Stmt Lower(const LowerArgs &T, arith::Analyzer *analyzer) const;
  LayoutMap InferLayout(const LayoutInferArgs &T, InferLevel level) const;
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
  enum class GemmWarpPolicy {
    kSquare = 0,
    kFullRow = 1,
    kFullCol = 2,
  } policy;

  std::pair<int, int>
  ComputeWarpPartition(int num_warps, Target target,
                       bool maybe_hopper_wgmma = true) const;

  Array<PrimExpr> call_args;
  tir::Buffer A, B, C, E;
  bool trans_A, trans_B;
  int M, N, K;
  bool clear_accum = false;
  // k_pack please ref to bitblas/tl/mfma_macro_generator.py::k_pack
  // only will be enabled under cdna mfma instructions
  int kPack = 1;
  int wg_wait = 0;
94
95
96
97
98
99
100
101
102
103
104
105

  TileOperator Clone() const;

private:
  mutable bool completed_ = false;
};

class GemmSP : public TileOperator {
public:
  TVM_DEFINE_OBJECT_REF_METHODS(GemmSP, TileOperator, GemmSPNode);
  TVM_DLL GemmSP(Array<PrimExpr> args, BufferMap vmap);
  static const Op &Get();
106
107
108
109
110
111
};

} // namespace tl
} // namespace tvm

#endif //  TVM_TL_OP_GEMM_SP_H_