Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
33 commits
Select commit Hold shift + click to select a range
cd11fcb
reimplement ThreadSync pass
silentCoder-dev Jan 7, 2026
fcd3659
fix bugs
silentCoder-dev Jan 7, 2026
fab1648
remove old storage_access
silentCoder-dev Jan 7, 2026
bd73f7a
remove deadcode for debugging
silentCoder-dev Jan 7, 2026
fea9123
format
silentCoder-dev Jan 7, 2026
e3541a6
add WhileOp & expose MakeGuard in constr_visitor
silentCoder-dev Jan 8, 2026
bf85fcb
bugfix
silentCoder-dev Jan 8, 2026
5462282
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
silentCoder-dev Jan 8, 2026
5918b7b
bugfix
silentCoder-dev Jan 8, 2026
e1fdb59
bugfix
silentCoder-dev Jan 8, 2026
eeede46
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
silentCoder-dev Jan 8, 2026
7d91bca
bugfix
silentCoder-dev Jan 9, 2026
d05b15e
trigger ci
silentCoder-dev Jan 9, 2026
f9f8e84
add try for checking
silentCoder-dev Jan 9, 2026
00f0fae
bugfix
silentCoder-dev Jan 9, 2026
1aa8bd4
update tvm with better analyzer
silentCoder-dev Jan 9, 2026
c8e3f6b
update tvm
silentCoder-dev Jan 12, 2026
9c5341b
Log for debugging in thread_storage_sync
silentCoder-dev Jan 12, 2026
00db1ff
format
silentCoder-dev Jan 12, 2026
d61093d
consider let node
silentCoder-dev Jan 12, 2026
1b5ddef
typo
silentCoder-dev Jan 12, 2026
a8de275
remove debugging log
silentCoder-dev Jan 12, 2026
35a8c8e
consider while op in thread_sync & fix bug sbout letnode
silentCoder-dev Jan 12, 2026
53fc9ef
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
silentCoder-dev Jan 12, 2026
2d7db7d
add testing for thread_sync issues
silentCoder-dev Jan 12, 2026
c3991e9
update testing for thread_sync
silentCoder-dev Jan 12, 2026
6353023
format
silentCoder-dev Jan 12, 2026
4c9bd2b
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
silentCoder-dev Jan 13, 2026
f3d377f
consider T.Ramp expr
silentCoder-dev Jan 13, 2026
12839ad
remove useless header
silentCoder-dev Jan 13, 2026
4b41d4c
format
silentCoder-dev Jan 13, 2026
2dd4308
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
LeiWang1999 Jan 13, 2026
e60b339
Merge branch 'main' of https://github.com/tile-ai/tilelang into threa…
LeiWang1999 Jan 14, 2026
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
51 changes: 51 additions & 0 deletions src/transform/common/constr_visitor.h
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@
#include "tvm/tir/op.h"
#include "tvm/tir/stmt.h"
#include "tvm/tir/var.h"
#include <ostream>
#include <tvm/ffi/reflection/registry.h>
#include <tvm/tir/analysis.h>
#include <tvm/tir/builtin.h>
Expand Down Expand Up @@ -40,6 +41,31 @@ struct Constr {
Constr(Constr &&other) = default;
Constr &operator=(const Constr &other) = default;

void format(std::ostream &os) const {
os << "Constr(kind=";
switch (kind) {
case kConstr:
os << "kConstr";
os << ", is_assume=" << (is_assume ? "true" : "false");
os << ", value=" << value;
break;
case kBindValue:
os << "kBindValue";
os << ", var=" << var->name_hint;
os << ", value=" << value;
break;
case kBindRange:
os << "kBindRange";
os << ", var=" << var->name_hint;
os << ", range=Range(min=" << range->min;
os << ", extent=" << range->extent << ")";
break;
default:
os << "Unknown";
}
os << ")";
}

PrimExpr ToGenericConstr() const {
switch (kind) {
case kConstr:
Expand Down Expand Up @@ -98,22 +124,38 @@ struct ConstrSet {
constrs_.push_back(c);
}
}

void format(std::ostream &os) const {
os << "ConstrSet(size=" << constrs_.size() << ") {\n";
for (size_t i = 0; i < constrs_.size(); ++i) {
os << " [" << i << "] ";
constrs_[i].format(os);
os << "\n";
}
os << "}";
}

std::vector<Constr> constrs_;
};

struct ConstrVisitor : public tir::StmtExprVisitor {
private:
using Base = tir::StmtExprVisitor;

struct Guard {
std::vector<Constr> &constrs;
~Guard() { constrs.pop_back(); }
};

protected:
template <typename... Args> Guard MakeGuard(const Args... args) {
constr_stack_.push_back(Constr(args...));
return Guard{constr_stack_};
}

public:
using StmtExprVisitor::VisitExpr_;
using StmtExprVisitor::VisitStmt_;
void VisitIfThenElseExpr(const PrimExpr cond, const PrimExpr true_value,
const PrimExpr false_value) {
{
Expand Down Expand Up @@ -178,9 +220,18 @@ struct ConstrVisitor : public tir::StmtExprVisitor {
auto guard_2 = MakeGuard(op->extent > 0);
Base::VisitStmt_(op);
} else {
auto guard_1 =
MakeGuard(op->loop_var, Range::FromMinExtent(op->min, op->extent));
auto guard_2 = MakeGuard(op->extent > 0);
Base::VisitStmt_(op);
}
}
void VisitStmt_(const tir::WhileNode *op) override {
{
auto guard = MakeGuard(op->condition);
Base::VisitStmt(op->body);
}
}
std::vector<Constr> constr_stack_;
};
} // namespace tvm::tl
Expand Down
1 change: 0 additions & 1 deletion src/transform/eliminate_storage_sync_for_mbarrier.cc
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,6 @@
* \file eliminate_storage_sync_for_mbarrier.cc
*/
#include "../op/builtin.h"
#include "./storage_access.h"
#include "arith/ir_mutator_with_analyzer.h"
#include "arith/ir_visitor_with_analyzer.h"
#include <tvm/ffi/function.h>
Expand Down
1 change: 0 additions & 1 deletion src/transform/inject_ptx_async_copy.cc
Original file line number Diff line number Diff line change
Expand Up @@ -29,7 +29,6 @@
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/transform.h>

#include "storage_access.h"
#include "tir/ir/buffer_common.h"
#include "tvm/tir/stmt.h"

Expand Down
Loading
Loading