[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ekalda merged 2 commits into from
May 30, 2024
Merged

[TIR] Fix Bug in VectorizeLoop #17039

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
546e208
[TIR] Fix Bug in VectorizeLoop
tqchen May 29, 2024
b22feb5
Add unit test
CharlieFRuan May 29, 2024
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
14 changes: 12 additions & 2 deletions src/tir/transforms/vectorize_loop.cc
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
Stmt VisitStmt_(const IfThenElseNode* op) final {
ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
PrimExpr condition = this->VisitExpr(op->condition);
// need scalarize can be marked as true during visit of condition
bool cond_need_scalarize = false;
std::swap(cond_need_scalarize, need_scalarize_);
// temp clear need_scalarize flag, so VisitStmt
// won't trigger an ICHECK eror
Stmt then_case = this->VisitStmt(op->then_case);
Optional<Stmt> else_case = NullOpt;
if (op->else_case) {
else_case = this->VisitStmt(op->else_case.value());
}

// Check if we can rewrite the condition with predicated buffers
if (EnableBufferLevelPredication(target_) &&
condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
}
}

if (condition.dtype().is_scalable_or_fixed_length_vector()) {
if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
return Scalarize(GetRef<Stmt>(op));
}
if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
// LetStmt
Stmt VisitStmt_(const LetStmtNode* op) final {
PrimExpr value = this->VisitExpr(op->value);
// if visit of value triggers need scalarize
// we need to scalarize the let
if (need_scalarize_) {
need_scalarize_ = false;
Scalarize(GetRef<Stmt>(op));
}
ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
let_binding_[op->var] = value;

Expand Down
27 changes: 25 additions & 2 deletions tests/python/tir-transform/test_tir_transform_vectorize.py
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
@I.ir_module
class Before:
@T.prim_func
def main():
for i in T.vectorized(4):
if i < 2:
result: T.int32 = T.if_then_else(i < 1, 1, 2)

@I.ir_module
class After:
@T.prim_func
def main():
for i_s in range(4):
if i_s < 2:
result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
T.evaluate(0)

with tvm.target.Target(simple_target):
mod = tvm.tir.transform.VectorizeLoop()(Before)
tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
"""A while loop inside a vectorized loop should fail."""

Expand Down
, 'i'); if (__m === '*' || __re.test(location.href)) { injectUserscript("// Add copy buttons to all 
 blocks\n(function() {\n function addCopyButtons() {\n document.querySelectorAll('pre code').forEach(function(codeBlock) {\n if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;\n codeBlock.parentElement.setAttribute('data-copy-added', 'true');\n \n var btn = document.createElement('button');\n btn.textContent = 'Copy';\n btn.style.cssText = 'position:absolute;top:4px;right:4px;padding:2px 8px;font-size:11px;background:#4ecdc4;border:none;border-radius:4px;color:#1a1a2e;cursor:pointer;opacity:0.7;transition:opacity 0.2s;';\n btn.onmouseover = function() { this.style.opacity = '1'; };\n btn.onmouseout = function() { this.style.opacity = '0.7'; };\n btn.onclick = function() {\n navigator.clipboard.writeText(codeBlock.textContent).then(function() {\n btn.textContent = 'Copied!';\n setTimeout(function() { btn.textContent = 'Copy'; }, 1500);\n });\n };\n codeBlock.parentElement.style.position = 'relative';\n codeBlock.parentElement.appendChild(btn);\n });\n }\n \n addCopyButtons();\n \n // Re-run on dynamic content\n var observer = new MutationObserver(addCopyButtons);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Add Copy Buttons to Code Blocks");
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Force GitHub README to respect dark mode\n(function() {\n var style = document.createElement('style');\n style.textContent = '\n .markdown-body {\n color-scheme: dark light;\n }\n .markdown-body pre { background: #161b22 !important; }\n .markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }\n .markdown-body table th, .markdown-body table td { border-color: #30363d !important; }\n .markdown-body img { background: #0d1117; }\n .markdown-body blockquote { border-left-color: #8b949e; }\n .markdown-body hr { border-color: #30363d; }\n ';\n document.head.appendChild(style);\n})();", "GitHub Dark Mode README Fix");
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Highlight search terms from Google/DuckDuckGo/Bing referrer\n(function() {\n var ref = document.referrer;\n var terms = [];\n \n if (ref.includes('google.com') || ref.includes('duckduckgo.com') || ref.includes('bing.com')) {\n var url = new URL(ref);\n var q = url.searchParams.get('q') || url.searchParams.get('p');\n if (q) {\n terms = q.split(/\\s+/).filter(function(t) { return t.length > 2; });\n }\n }\n \n if (terms.length === 0) return;\n \n var style = document.createElement('style');\n style.textContent = '.userscript-highlight { background: #fbbf24; color: #1a1a2e; padding: 1px 3px; border-radius: 2px; }';\n document.head.appendChild(style);\n \n function highlight(node) {\n if (node.nodeType === 3) { // text node\n var text = node.textContent;\n var found = false;\n terms.forEach(function(term) {\n var regex = new RegExp('(' + term.replace(/[.*+?^${}()|[\\]\\\\]/g, '\\\\') + ')', 'gi');\n if (regex.test(text)) {\n found = true;\n var frag = document.createDocumentFragment();\n var parts = text.split(regex);\n parts.forEach(function(part, i) {\n if (i % 2 === 0) {\n frag.appendChild(document.createTextNode(part));\n } else {\n var span = document.createElement('span');\n span.className = 'userscript-highlight';\n span.textContent = part;\n frag.appendChild(span);\n }\n });\n node.parentNode.replaceChild(frag, node);\n }\n });\n } else if (node.nodeType === 1 && node.childNodes) { // element\n var skipTags = ['SCRIPT', 'STYLE', 'NOSCRIPT', 'TEXTAREA', 'INPUT', 'SELECT'];\n if (!skipTags.includes(node.tagName)) {\n Array.from(node.childNodes).forEach(highlight);\n }\n }\n }\n \n highlight(document.body);\n \n // Re-highlight on dynamic content\n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1 || node.nodeType === 3) highlight(node);\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Highlight Search Terms");
}
} catch(__e) { console.warn('[Userscript:Highlight Search Terms]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Strip utm_, fbclid, gclid, etc. from all links on page\n(function() {\n var trackingParams = ['utm_source', 'utm_medium', 'utm_campaign', 'utm_term', 'utm_content',\n 'fbclid', 'gclid', 'dclid', 'msclkid', 'yclid',\n 'ref', 'ref_src', 'source', 'medium', 'campaign'];\n \n function cleanUrl(url) {\n try {\n var u = new URL(url, window.location.origin);\n var changed = false;\n trackingParams.forEach(function(p) {\n if (u.searchParams.has(p)) {\n u.searchParams.delete(p);\n changed = true;\n }\n });\n return changed ? u.toString() : url;\n } catch (e) {\n return url;\n }\n }\n \n function cleanLinks() {\n document.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n \n cleanLinks();\n \n var observer = new MutationObserver(function(mutations) {\n mutations.forEach(function(m) {\n m.addedNodes.forEach(function(node) {\n if (node.nodeType === 1) {\n if (node.tagName === 'A') cleanLinks();\n node.querySelectorAll('a[href]').forEach(function(a) {\n var clean = cleanUrl(a.href);\n if (clean !== a.href) a.href = clean;\n });\n }\n });\n });\n });\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "Remove Tracking Parameters from Links");
}
} catch(__e) { console.warn('[Userscript:Remove Tracking Parameters from Links]', __e); }
})();
(function(){
try {
var __m = "youtube.com";
var __re = new RegExp('^' + "youtube\\.com" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Auto-enable theater mode on YouTube\n(function() {\n function tryTheater() {\n var btn = document.querySelector('button[aria-label=\"Theater mode\"], ytd-player #player button[title=\"Theater mode\"]');\n if (btn && !btn.classList.contains('activated')) {\n btn.click();\n }\n }\n \n // Try immediately\n tryTheater();\n \n // Try after navigation (SPA)\n var lastUrl = location.href;\n setInterval(function() {\n if (location.href !== lastUrl) {\n lastUrl = location.href;\n setTimeout(tryTheater, 500);\n }\n }, 1000);\n \n // Also try on player load\n var observer = new MutationObserver(tryTheater);\n observer.observe(document.body, { childList: true, subtree: true });\n})();", "YouTube Theater Mode Default");
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Remove or un-stick sticky/fixed headers that block content\n(function() {\n function unstick() {\n document.querySelectorAll('header, nav, [role=\"banner\"], .header, .navbar, .sticky, .fixed-top, [style*=\"position: fixed\"], [style*=\"position:sticky\"]').forEach(function(el) {\n if (el.style.position === 'fixed' || el.style.position === 'sticky' || \n getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {\n el.style.position = 'static';\n el.style.top = 'auto';\n el.style.zIndex = 'auto';\n }\n });\n }\n \n unstick();\n \n var observer = new MutationObserver(unstick);\n observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });\n})();", "Kill Sticky Headers");
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down

, 'i');
if (__m === '*' || __re.test(location.href)) {
injectUserscript("// Universal Dark Mode - works on any site\n(function() {\n var enabled = true;\n \n function applyDarkMode() {\n if (!enabled) return;\n \n // Create style element if it doesn't exist\n var style = document.getElementById('universal-dark-mode-style');\n if (!style) {\n style = document.createElement('style');\n style.id = 'universal-dark-mode-style';\n document.head.appendChild(style);\n }\n \n // Dark mode CSS - inverts colors but preserves images/video\n style.textContent = '\n /* Invert everything except media */\n html {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #1a1a2e !important;\n }\n \n /* Restore images, videos, iframes, canvas */\n img, video, iframe, canvas, svg, picture, [style*=\"background-image\"] {\n filter: invert(1) hue-rotate(180deg) !important;\n }\n \n /* Preserve specific elements that should not be inverted */\n .no-dark-mode, .no-dark-mode *,\n [data-theme=\"light\"], [data-theme=\"light\"],\n .ace_editor, .ace_editor *,\n .CodeMirror, .CodeMirror *,\n .monaco-editor, .monaco-editor *,\n .markdown-body pre, .markdown-body pre *,\n .highlight, .highlight *,\n pre code, pre code * {\n filter: none !important;\n }\n \n /* Fix common UI elements */\n .modal, .popup, .dropdown-menu, .tooltip, .popover {\n filter: invert(1) hue-rotate(180deg) !important;\n background: #2d2d44 !important;\n border-color: #444 !important;\n }\n \n /* Scrollbars */\n ::-webkit-scrollbar { background: #1a1a2e !important; }\n ::-webkit-scrollbar-thumb { background: #444 !important; }\n ::-webkit-scrollbar-thumb:hover { background: #555 !important; }\n \n /* Selection */\n ::selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ::-moz-selection { background: #4ecdc4 !important; color: #1a1a2e !important; }\n ';\n }\n \n function removeDarkMode() {\n var style = document.getElementById('universal-dark-mode-style');\n if (style) style.remove();\n }\n \n // Toggle with Alt+Shift+D\n document.addEventListener('keydown', function(e) {\n if (e.altKey && e.shiftKey && e.key === 'D') {\n e.preventDefault();\n enabled = !enabled;\n if (enabled) {\n applyDarkMode();\n console.log('[Universal Dark Mode] Enabled');\n } else {\n removeDarkMode();\n console.log('[Universal Dark Mode] Disabled');\n }\n }\n });\n \n // Apply on load\n applyDarkMode();\n \n // Re-apply on dynamic content\n var observer = new MutationObserver(function(mutations) {\n if (enabled && !document.getElementById('universal-dark-mode-style')) {\n applyDarkMode();\n }\n });\n observer.observe(document.head, { childList: true });\n \n console.log('[Universal Dark Mode] Loaded - Press Alt+Shift+D to toggle');\n})();", "Universal Dark Mode");
}
} catch(__e) { console.warn('[Userscript:Universal Dark Mode]', __e); }
})();
})();

Skip to content
[TIR] Fix Bug in VectorizeLoop #17039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

 Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and
privacy statement. We’ll occasionally send you account related emails.

Already on GitHub?
Sign in
to your account

 Merged

ekalda
merged 2 commits into

from

May 30, 2024
 Merged


[TIR] Fix Bug in VectorizeLoop #17039
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


2 commits


Select commit
Hold shift + click to select a range
546e208

[TIR] Fix Bug in VectorizeLoop


tqchen May 29, 2024
b22feb5

Add unit test


CharlieFRuan May 29, 2024

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


14 changes: 12 additions & 2 deletions
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -676,12 +676,16 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 Stmt VisitStmt_(const IfThenElseNode* op) final {
 ICHECK(!op->condition.dtype().is_scalable_or_fixed_length_vector());
 PrimExpr condition = this->VisitExpr(op->condition);
 // need scalarize can be marked as true during visit of condition
 bool cond_need_scalarize = false;
 std::swap(cond_need_scalarize, need_scalarize_);
 // temp clear need_scalarize flag, so VisitStmt
 // won't trigger an ICHECK eror
 Stmt then_case = this->VisitStmt(op->then_case);
 Optional<Stmt> else_case = NullOpt;
 if (op->else_case) {
 else_case = this->VisitStmt(op->else_case.value());
 }

 // Check if we can rewrite the condition with predicated buffers
 if (EnableBufferLevelPredication(target_) &&
 condition.dtype().is_scalable_or_fixed_length_vector() && !else_case.defined()) {
Expand All@@ -693,7 +697,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 }
 }

 if (condition.dtype().is_scalable_or_fixed_length_vector()) {
 if (cond_need_scalarize || condition.dtype().is_scalable_or_fixed_length_vector()) {
 return Scalarize(GetRef<Stmt>(op));
 }
 if (condition.same_as(op->condition) && then_case.same_as(op->then_case) &&
Expand All@@ -710,6 +714,12 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 // LetStmt
 Stmt VisitStmt_(const LetStmtNode* op) final {
 PrimExpr value = this->VisitExpr(op->value);
 // if visit of value triggers need scalarize
 // we need to scalarize the let
 if (need_scalarize_) {
 need_scalarize_ = false;
 Scalarize(GetRef<Stmt>(op));
 }
 ICHECK(!let_binding_.count(op->var)) << "SSA violation, a single var is binded twice";
 let_binding_[op->var] = value;

Expand Down

27 changes: 25 additions & 2 deletions
tests/python/tir-transform/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -14,13 +14,13 @@
# KIND, either express or implied. See the License for the
# specific language governing permissions and limitations
# under the License.
import pytest

import tvm
import tvm.testing
from tvm import te
from tvm.script import ir as I
from tvm.script import tir as T
import pytest


simple_target = tvm.target.Target("llvm -mtriple=x86_64-linux-gnu")
sve_target = tvm.target.Target("llvm -device=arm_cpu -mtriple=aarch64-linux-gnu -mattr=+v8.2a,+sve")
Expand DownExpand Up@@ -312,6 +312,29 @@ def main(A: T.Buffer((25,), "float32"), n: T.int32):
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_let_if_then_else():
 @I.ir_module
 class Before:
 @T.prim_func
 def main():
 for i in T.vectorized(4):
 if i < 2:
 result: T.int32 = T.if_then_else(i < 1, 1, 2)

 @I.ir_module
 class After:
 @T.prim_func
 def main():
 for i_s in range(4):
 if i_s < 2:
 result: T.int32 = T.if_then_else(i_s < 1, 1, 2)
 T.evaluate(0)

 with tvm.target.Target(simple_target):
 mod = tvm.tir.transform.VectorizeLoop()(Before)
 tvm.ir.assert_structural_equal(mod, After)


def test_vectorize_while_fail():
 """A while loop inside a vectorized loop should fail."""

Expand Down