You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/09/16 18:47:10 UTC
[GitHub] [tvm] MasterJH5574 opened a new pull request, #12819: [BugFix][TIR] Fix Buffer LCA Detector
MasterJH5574 opened a new pull request, #12819:
URL: https://github.com/apache/tvm/pull/12819
Prior to this PR, the LCA detector of buffers in TIR didn't take buffer memory scopes and GPU hierarchy into consideration. An consequent issue is that, when an intermediate buffer is in global memory, TIR's lowering passes don't necessarily allocated the intermediate buffer outside all `blockIdx`. As a result, the global intermediate buffer is allocated under a GPU thread block, which is illegal.
This PR fixes this issue by fixing the LCA detector, making it be aware of the buffer memory scopes and GPU hierarchy. With this fix, the global intermediate buffers are all allocated outside `blockIdx`.
cc @junrushao @tqchen @Hzfengsy @spectrometerHBH
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tqchen commented on a diff in pull request #12819: [BugFix][TIR] Fix Buffer LCA Detector
Posted by GitBox <gi...@apache.org>.
tqchen commented on code in PR #12819:
URL: https://github.com/apache/tvm/pull/12819#discussion_r973330490
##########
src/tir/analysis/buffer_access_lca_detector.cc:
##########
@@ -150,6 +176,17 @@ class LCADetector : public StmtExprVisitor {
}
}
+ void UpdateWithBlockidx() {
+ for (const auto& it : buffer_lca_) {
+ if (GetRef<Buffer>(it.first).scope() == "global") {
Review Comment:
Use StorageScope here https://github.com/apache/tvm/blob/main/src/runtime/thread_storage_scope.h#L86
Consider move up to the ThreadScope as long as their rank is higher equal than threadScope.rank
##########
src/tir/analysis/buffer_access_lca_detector.cc:
##########
@@ -107,6 +122,17 @@ class LCADetector : public StmtExprVisitor {
ancestor_scopes_.pop_back();
}
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == attr::thread_extent) {
+ const auto* iter = op->node.as<IterVarNode>();
+ ICHECK_NOTNULL(iter);
+ if (StartsWith(iter->thread_tag, "blockIdx.")) {
Review Comment:
Instead of string matching, use ThreadScope here https://github.com/apache/tvm/blob/main/src/runtime/thread_storage_scope.h#L161
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] Hzfengsy merged pull request #12819: [BugFix][TIR] Fix Buffer LCA Detector
Posted by GitBox <gi...@apache.org>.
Hzfengsy merged PR #12819:
URL: https://github.com/apache/tvm/pull/12819
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] MasterJH5574 commented on a diff in pull request #12819: [BugFix][TIR] Fix Buffer LCA Detector
Posted by GitBox <gi...@apache.org>.
MasterJH5574 commented on code in PR #12819:
URL: https://github.com/apache/tvm/pull/12819#discussion_r973392746
##########
src/tir/analysis/buffer_access_lca_detector.cc:
##########
@@ -107,6 +122,17 @@ class LCADetector : public StmtExprVisitor {
ancestor_scopes_.pop_back();
}
+ void VisitStmt_(const AttrStmtNode* op) final {
+ if (op->attr_key == attr::thread_extent) {
+ const auto* iter = op->node.as<IterVarNode>();
+ ICHECK_NOTNULL(iter);
+ if (StartsWith(iter->thread_tag, "blockIdx.")) {
Review Comment:
Thanks for pointing it out!! Just updated :-)
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org