gemm_sp.h 3.05 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 "gemm.h"
11
#include "operator.h"
12
13

namespace tvm {
14

15
16
17
18
namespace tl {

using namespace tir;

19
20
class GemmSPWarpPolicyNode : public GemmWarpPolicyNode {
public:
21
  std::pair<int, int> computeWarpPartition(int M, int N, int block_size,
22
23
                                           Target target, bool use_wgmma,
                                           int bits) const;
24
25
  TVM_FFI_DECLARE_OBJECT_INFO("tl.GemmSPWarpPolicy", GemmSPWarpPolicyNode,
                              GemmWarpPolicyNode);
26
27
28
29
};

class GemmSPWarpPolicy : public ObjectRef {
public:
30
31
  TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(GemmSPWarpPolicy, ObjectRef,
                                             GemmSPWarpPolicyNode);
32
33

  explicit GemmSPWarpPolicy(GemmWarpPolicyType policy_type) {
34
    auto node = tvm::ffi::make_object<GemmSPWarpPolicyNode>();
35
36
37
38
39
    node->policy_type = (int)policy_type;
    data_ = std::move(node);
  }

  explicit GemmSPWarpPolicy(int policy_type) {
40
    auto node = tvm::ffi::make_object<GemmSPWarpPolicyNode>();
41
42
43
44
45
    node->policy_type = policy_type;
    data_ = std::move(node);
  }

  explicit GemmSPWarpPolicy(int m_warp, int n_warp) {
46
    auto node = tvm::ffi::make_object<GemmSPWarpPolicyNode>();
47
48
49
50
51
52
53
    node->m_warp = m_warp;
    node->n_warp = n_warp;
    node->policy_type = (int)GemmWarpPolicyType::kFree;
    data_ = std::move(node);
  }
};

54
class GemmSPNode : public TileOperatorNode {
55
public:
56
57
58
59
  tir::Buffer a_, b_, c_, e_;
  bool transA_, transB_;
  int m_, n_, k_;
  bool clearAccum_ = false;
60
61
  // k_pack please ref to bitblas/tl/mfma_macro_generator.py::k_pack
  // only will be enabled under cdna mfma instructions
62
63
  int kPack_ = 1;
  int wgWait_ = 0;
64

65
  mutable GemmSPWarpPolicy policy_;
66

67
  TVM_FFI_DECLARE_OBJECT_INFO_FINAL("tl.GemmSP", GemmSPNode, TileOperatorNode);
68
69
70
71
  Stmt Lower(const LowerArgs &T, arith::Analyzer *analyzer) const override;
  LayoutMap InferLayout(const LayoutInferArgs &T,
                        InferLevel level) const override;

72
73
  TileOperator Clone() const;

74
75
76
  static void RegisterReflection() {
    namespace refl = tvm::ffi::reflection;
    refl::ObjectDef<GemmSPNode>()
77
78
79
80
81
82
83
84
85
86
87
88
89
        .def_ro("policy", &GemmSPNode::policy_)
        .def_ro("a", &GemmSPNode::a_)
        .def_ro("b", &GemmSPNode::b_)
        .def_ro("c", &GemmSPNode::c_)
        .def_ro("e", &GemmSPNode::e_)
        .def_ro("transA", &GemmSPNode::transA_)
        .def_ro("transB", &GemmSPNode::transB_)
        .def_ro("m", &GemmSPNode::m_)
        .def_ro("n", &GemmSPNode::n_)
        .def_ro("k", &GemmSPNode::k_)
        .def_ro("clearAccum", &GemmSPNode::clearAccum_)
        .def_ro("kPack", &GemmSPNode::kPack_)
        .def_ro("wgWait", &GemmSPNode::wgWait_);
90
91
  }

92
93
94
95
96
97
private:
  mutable bool completed_ = false;
};

class GemmSP : public TileOperator {
public:
98
  TVM_FFI_DEFINE_OBJECT_REF_METHODS_NULLABLE(GemmSP, TileOperator, GemmSPNode);
99
100
  TVM_DLL GemmSP(Array<PrimExpr> args, BufferMap vmap);
  static const Op &Get();
101
102
103
104
105
106
};

} // namespace tl
} // namespace tvm

#endif //  TVM_TL_OP_GEMM_SP_H_