Skip to content
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
181 changes: 181 additions & 0 deletions src/transform/common/constr_visitor.h
Original file line number Diff line number Diff line change
@@ -0,0 +1,181 @@
#include "tvm/arith/analyzer.h"
#include "tvm/ffi/base_details.h"
#include "tvm/ffi/object.h"
#include "tvm/ir/expr.h"
#include "tvm/tir/op.h"
#include "tvm/tir/stmt.h"
#include "tvm/tir/var.h"
#include <tvm/ffi/reflection/registry.h>
#include <tvm/tir/analysis.h>
#include <tvm/tir/builtin.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/transform.h>
Comment thread
coderabbitai[bot] marked this conversation as resolved.

Comment thread
kurisu6912 marked this conversation as resolved.
namespace tvm::tl {

struct Constr {

enum Kind {
kConstr,
kBindValue,
kBindRange,
} kind;
bool is_assume = false;
tir::Var var;
PrimExpr value;
Range range;

Constr(PrimExpr constr, bool is_assume = false)
: kind(kConstr), value(constr), is_assume(is_assume) {};
Constr(tir::Var var, PrimExpr val)
: kind(kBindValue), var(var), value(val) {};
Constr(tir::Var var, Range range)
: kind(kBindRange), var(var), range(range) {};

Constr() = default;
Constr(const Constr &other) = default;
Constr(Constr &&other) = default;
Constr &operator=(const Constr &other) = default;

PrimExpr ToGenericConstr() const {
switch (kind) {
case kConstr:
return value;
case kBindValue:
return var == value;
case kBindRange:
return tir::And(var >= range->min, var < (range->min + range->extent));
}
LOG(FATAL) << "Unreachable";
return PrimExpr();
}
Constr Substitute(ffi::Map<tir::Var, PrimExpr> subs) const {
return Constr(tir::Substitute(ToGenericConstr(), subs));
}
void Populate(arith::Analyzer &analyzer) const {
switch (kind) {
case kConstr:
analyzer.EnterConstraint(value);
break;
case kBindValue:
analyzer.Bind(var, value);
break;
case kBindRange:
analyzer.Bind(var, range);
break;
}
}
Comment thread
coderabbitai[bot] marked this conversation as resolved.
};

struct ConstrSet {
ConstrSet Substitute(ffi::Map<tir::Var, PrimExpr> subs) const {
ConstrSet new_set;
for (const auto &c : constrs_) {
new_set.constrs_.push_back(c.Substitute(subs));
}
return new_set;
}
void Populate(arith::Analyzer &analyzer) const {
for (const auto &c : constrs_) {
c.Populate(analyzer);
}
}
bool CanProve(const PrimExpr &expr) const {
arith::Analyzer analyzer;
Populate(analyzer);
return analyzer.CanProve(expr);
}
template <typename... Args> void AddConstr(Args... args) {
constrs_.push_back(Constr(args...));
}
void Extend(const ConstrSet &other) {
for (const auto &c : other.constrs_) {
constrs_.push_back(c);
}
}
std::vector<Constr> constrs_;
};

struct ConstrVisitor : public tir::StmtExprVisitor {
private:
using Base = tir::StmtExprVisitor;
struct Guard {
std::vector<Constr> &constrs;
~Guard() { constrs.pop_back(); }
};
template <typename... Args> Guard MakeGuard(const Args... args) {
constr_stack_.push_back(Constr(args...));
return Guard{constr_stack_};
}

public:
void VisitIfThenElseExpr(const PrimExpr cond, const PrimExpr true_value,
const PrimExpr false_value) {
{
auto guard = MakeGuard(cond);
Base::VisitExpr(true_value);
}
{
auto guard = MakeGuard(tir::Not(cond));
Base::VisitExpr(false_value);
}
}
void VisitStmt_(const tir::LetStmtNode *op) override {
auto guard = MakeGuard(op->var, op->value);
Base::VisitStmt_(op);
}
void VisitStmt_(const tir::AttrStmtNode *op) override {
if (op->attr_key == tir::attr::tilelang_assume) {
auto expr = Downcast<PrimExpr>(op->node);
auto guard = MakeGuard(expr, true);
Base::VisitStmt_(op);
} else if (op->attr_key == tir::attr::thread_extent ||
op->attr_key == tir::attr::virtual_thread) {
tir::IterVar iv = Downcast<tir::IterVar>(op->node);
Range dom =
Range::FromMinExtent(tir::make_zero(op->value.dtype()), op->value);
auto guard = MakeGuard(iv->var, dom);
Base::VisitStmt_(op);
} else {
Base::VisitStmt_(op);
}
}
void VisitStmt_(const tir::AssertStmtNode *op) override {
auto guard = MakeGuard(op->condition);
Base::VisitStmt_(op);
}
void VisitStmt_(const tir::IfThenElseNode *op) override {
{
auto guard = MakeGuard(op->condition);
Base::VisitStmt(op->then_case);
}
if (op->else_case) {
auto guard = MakeGuard(tir::Not(op->condition));
Base::VisitStmt(op->else_case.value());
}
}
void VisitExpr_(const tir::SelectNode *op) override {
VisitIfThenElseExpr(op->condition, op->true_value, op->false_value);
}
void VisitExpr_(const tir::CallNode *op) override {
static auto op_if_then_else = Op::Get("tir.if_then_else");
if (op->op.same_as(op_if_then_else)) {
VisitIfThenElseExpr(op->args[0], op->args[1], op->args[2]);
} else {
Base::VisitExpr_(op);
}
}
void VisitStmt_(const tir::ForNode *op) override {
if (op->kind == tir::ForKind::kParallel ||
op->kind == tir::ForKind::kVectorized) {
auto guard_1 =
MakeGuard(op->loop_var, Range::FromMinExtent(op->min, op->extent));
auto guard_2 = MakeGuard(op->extent > 0);
Base::VisitStmt_(op);
} else {
Base::VisitStmt_(op);
}
}
std::vector<Constr> constr_stack_;
};
} // namespace tvm::tl
165 changes: 3 additions & 162 deletions src/transform/verify_parallel_loop.cc
Original file line number Diff line number Diff line change
@@ -1,3 +1,4 @@
#include "common/constr_visitor.h"
#include "layout_reducer.h"
#include "tvm/arith/analyzer.h"
#include "tvm/ffi/base_details.h"
Expand All @@ -17,168 +18,8 @@ namespace tvm::tl {
using namespace tir;

namespace {

struct Constr {

enum Kind {
kConstr,
kBindValue,
kBindRange,
} kind;
bool is_assume = false;
Var var;
PrimExpr value;
Range range;

Constr(PrimExpr constr, bool is_assume = false)
: kind(kConstr), value(constr), is_assume(is_assume) {};
Constr(Var var, PrimExpr val) : kind(kBindValue), var(var), value(val) {};
Constr(Var var, Range range) : kind(kBindRange), var(var), range(range) {};

Constr() = default;
Constr(const Constr &other) = default;
Constr(Constr &&other) = default;
Constr &operator=(const Constr &other) = default;

PrimExpr ToGenericConstr() const {
switch (kind) {
case kConstr:
return value;
case kBindValue:
return var == value;
case kBindRange:
return And(var >= range->min, var < (range->min + range->extent));
}
LOG(FATAL) << "Unreachable";
return PrimExpr();
}
Constr Substitute(ffi::Map<Var, PrimExpr> subs) const {
return Constr(tir::Substitute(ToGenericConstr(), subs));
}
void Populate(arith::Analyzer &analyzer) const {
switch (kind) {
case kConstr:
analyzer.EnterConstraint(value);
break;
case kBindValue:
analyzer.Bind(var, value);
break;
case kBindRange:
analyzer.Bind(var, range);
break;
}
}
};

struct ConstrSet {
ConstrSet Substitute(ffi::Map<Var, PrimExpr> subs) const {
ConstrSet new_set;
for (const auto &c : constrs_) {
new_set.constrs_.push_back(c.Substitute(subs));
}
return new_set;
}
void Populate(arith::Analyzer &analyzer) const {
for (const auto &c : constrs_) {
c.Populate(analyzer);
}
}
bool CanProve(const PrimExpr &expr) const {
arith::Analyzer analyzer;
Populate(analyzer);
return analyzer.CanProve(expr);
}
template <typename... Args> void AddConstr(Args... args) {
constrs_.push_back(Constr(args...));
}
void Extend(const ConstrSet &other) {
for (const auto &c : other.constrs_) {
constrs_.push_back(c);
}
}
std::vector<Constr> constrs_;
};

struct ConstrVisitor : public StmtExprVisitor {
private:
using Base = StmtExprVisitor;
struct Guard {
std::vector<Constr> &constrs;
~Guard() { constrs.pop_back(); }
};
template <typename... Args> Guard MakeGuard(const Args... args) {
constr_stack_.push_back(Constr(args...));
return Guard{constr_stack_};
}

public:
void VisitIfThenElseExpr(const PrimExpr cond, const PrimExpr true_value,
const PrimExpr false_value) {
{
auto guard = MakeGuard(cond);
Base::VisitExpr(true_value);
}
{
auto guard = MakeGuard(Not(cond));
Base::VisitExpr(false_value);
}
}
void VisitStmt_(const LetStmtNode *op) override {
auto guard = MakeGuard(op->var, op->value);
Base::VisitStmt_(op);
}
void VisitStmt_(const AttrStmtNode *op) override {
if (op->attr_key == tir::attr::tilelang_assume) {
auto expr = Downcast<PrimExpr>(op->node);
auto guard = MakeGuard(expr, true);
Base::VisitStmt_(op);
} else if (op->attr_key == tir::attr::thread_extent ||
op->attr_key == tir::attr::virtual_thread) {
IterVar iv = Downcast<IterVar>(op->node);
Range dom = Range::FromMinExtent(make_zero(op->value.dtype()), op->value);
auto guard = MakeGuard(iv->var, dom);
Base::VisitStmt_(op);
} else {
Base::VisitStmt_(op);
}
}
void VisitStmt_(const AssertStmtNode *op) override {
auto guard = MakeGuard(op->condition);
Base::VisitStmt_(op);
}
void VisitStmt_(const IfThenElseNode *op) override {
{
auto guard = MakeGuard(op->condition);
Base::VisitStmt(op->then_case);
}
if (op->else_case) {
auto guard = MakeGuard(Not(op->condition));
Base::VisitStmt(op->else_case.value());
}
}
void VisitExpr_(const SelectNode *op) override {
VisitIfThenElseExpr(op->condition, op->true_value, op->false_value);
}
void VisitExpr_(const CallNode *op) override {
static auto op_if_then_else = Op::Get("tir.if_then_else");
if (op->op.same_as(op_if_then_else)) {
VisitIfThenElseExpr(op->args[0], op->args[1], op->args[2]);
} else {
Base::VisitExpr_(op);
}
}
void VisitStmt_(const ForNode *op) override {
if (op->kind == ForKind::kParallel || op->kind == ForKind::kVectorized) {
auto guard_1 =
MakeGuard(op->loop_var, Range::FromMinExtent(op->min, op->extent));
auto guard_2 = MakeGuard(op->extent > 0);
Base::VisitStmt_(op);
} else {
Base::VisitStmt_(op);
}
}
std::vector<Constr> constr_stack_;
};
using tvm::tl::ConstrSet;
using tvm::tl::ConstrVisitor;

struct ParallelLoopVerifier : public ConstrVisitor {
std::vector<Var> parallel_loop_vars_;
Expand Down
Loading