[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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
masahi merged 2 commits into from
Feb 22, 2022
Merged

[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes #10172

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
ffe6649
fix ramp in vectorize and narrow data type
lazycal Feb 4, 2022
8f6c1a3
Improve comments
lazycal Feb 4, 2022
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
17 changes: 17 additions & 0 deletions src/tir/transforms/narrow_datatype.cc
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
return StmtExprMutator::VisitExpr_(op);
}

PrimExpr VisitExpr_(const RampNode* op) final {
PrimExpr base = VisitExpr(op->base);
PrimExpr stride = VisitExpr(op->stride);
if (base.same_as(op->base) && stride.same_as(op->stride)) {
return GetRef<PrimExpr>(op);
} else {
if (base.dtype().is_int()) {
ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();
Comment on lines +262 to +263

@ganlerganlerFeb 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in

tvm/src/tir/ir/expr.cc

Line 705 in 22c488e

Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) {

Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?

@lazycallazycalFeb 7, 2022

Copy link
Copy Markdown
ContributorAuthor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.

int bits = std::max(base.dtype().bits(), stride.dtype().bits());
DataType dtype = base.dtype().with_bits(bits);
if (base.dtype() != dtype) base = cast(dtype, base);
if (stride.dtype() != dtype) stride = cast(dtype, stride);
}
return Ramp(base, stride, op->lanes);
}
}

PrimExpr VisitExpr_(const SizeVarNode* op) final {
if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
if (vmap_.find(op) == vmap_.end()) {
Expand Down
2 changes: 1 addition & 1 deletion src/tir/transforms/vectorize_loop.cc
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
using StmtMutator::operator();

Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
ramp_ = Ramp(0, 1, var_lanes);
ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
}

Stmt VisitStmt(const Stmt& stmt) final {
Expand Down
25 changes: 24 additions & 1 deletion tests/python/unittest/test_tir_transform_narrow_datatype.py
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
binds = {}
arg_list = []
for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

mod = schedule_to_module(sch, args)
mod = tvm.tir.transform.StorageFlatten(64)(mod)
if extra_passes:
for p in extra_passes:
mod = p(mod)
return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
)


def test_ramp_dtype_consistency():
"""
for (i :int64, (int64)0, (int64)4) {
A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
}
The infer result:
base: int64 -> int64 (since i is involved in another int64 expr)
stride: int64 -> int32

Thus ramp should still use int64 for both stride and base after rewrite.
"""
n = tvm.tir.IntImm("int64", 4)
m = tvm.tir.IntImm("int64", 2)
A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
s = te.create_schedule(A.op)
s[A].vectorize(A.op.axis[1])
lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
test_basic()
test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
test_slice()
test_relay_basic()
test_relay_take()
test_ramp_dtype_consistency()
9 changes: 9 additions & 0 deletions tests/python/unittest/test_tir_transform_vectorize.py
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
assert expected in error_msg


def test_vectorize_dtype_mismatch():
n = tvm.tir.IntImm("int64", 4)
A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
s = te.create_schedule(A.op)
s[A].vectorize(A.op.axis[0])
tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
test_vectorize_vector()
test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
test_vectorize_with_ge_cond()
test_vectorize_let()
test_vectorize_while_fail()
test_vectorize_dtype_mismatch()
, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()

, '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 Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes#10172

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

masahi
merged 2 commits into

from

Feb 22, 2022
 Merged


[TIR] Fix Ramp int32~64 mismatch in VectorizeLoop and NarrowDataType passes
#10172
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
ffe6649

fix ramp in vectorize and narrow data type


lazycal Feb 4, 2022
8f6c1a3

Improve comments


lazycal Feb 4, 2022

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


17 changes: 17 additions & 0 deletions
src/tir/transforms/narrow_datatype.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -253,6 +253,23 @@ class DataTypeRewriter : public StmtExprMutator {
 return StmtExprMutator::VisitExpr_(op);
 }

 PrimExpr VisitExpr_(const RampNode* op) final {
 PrimExpr base = VisitExpr(op->base);
 PrimExpr stride = VisitExpr(op->stride);
 if (base.same_as(op->base) && stride.same_as(op->stride)) {
 return GetRef<PrimExpr>(op);
 } else {
 if (base.dtype().is_int()) {
 ICHECK(stride.dtype().is_int()) << "Ramp base is int but stride is " << stride.dtype();

Comment on lines

+262
 to 
+263

@ganlerganlerFeb 7, 2022

Copy link


Copy Markdown

Contributor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can simply assume that base and stride should be of integer types. However, I also noticed that in
tvm/src/tir/ir/expr.cc

Line 705
in
22c488e
Ramp::Ramp(PrimExpr base, PrimExpr stride, int lanes, Span span) { 
Such assumptions are not checked. I am a bit curious if there will be, say base/stride in float types?
@lazycallazycalFeb 7, 2022

Copy link


Copy Markdown

ContributorAuthor
Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if they can be floats. So I added that conservative line of code that only acts on integers. If they can only be integer we should add that ICHECK you mention.
 int bits = std::max(base.dtype().bits(), stride.dtype().bits());
 DataType dtype = base.dtype().with_bits(bits);
 if (base.dtype() != dtype) base = cast(dtype, base);
 if (stride.dtype() != dtype) stride = cast(dtype, stride);
 }
 return Ramp(base, stride, op->lanes);
 }
 }

 PrimExpr VisitExpr_(const SizeVarNode* op) final {
 if (visitor_.vmap.find(op) != visitor_.vmap.end()) {
 if (vmap_.find(op) == vmap_.end()) {
Expand Down

2 changes: 1 addition & 1 deletion
src/tir/transforms/vectorize_loop.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -101,7 +101,7 @@ class Vectorizer : public StmtMutator, public ExprFunctor<PrimExpr(const PrimExp
 using StmtMutator::operator();

 Vectorizer(Var var, int var_lanes) : var_(var), var_lanes_(var_lanes) {
 ramp_ = Ramp(0, 1, var_lanes);
 ramp_ = Ramp(IntImm(var->dtype, 0), IntImm(var->dtype, 1), var_lanes);
 }

 Stmt VisitStmt(const Stmt& stmt) final {
Expand Down

25 changes: 24 additions & 1 deletion
tests/python/unittest/test_tir_transform_narrow_datatype.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -27,7 +27,7 @@ def lower_stmt(params, stmt, target_bits):
 return stmt


def lower_sch(sch, args, target_bits):
def lower_sch(sch, args, target_bits, extra_passes=None):
 binds = {}
 arg_list = []
 for x in args:
Expand All@@ -42,6 +42,9 @@ def lower_sch(sch, args, target_bits):

 mod = schedule_to_module(sch, args)
 mod = tvm.tir.transform.StorageFlatten(64)(mod)
 if extra_passes:
 for p in extra_passes:
 mod = p(mod)
 return tvm.tir.transform.NarrowDataType(target_bits)(mod)["main"].body


Expand DownExpand Up@@ -255,6 +258,25 @@ def check(shape, index, target_bits, target_dtype):
 )


def test_ramp_dtype_consistency():
 """
 for (i :int64, (int64)0, (int64)4) {
 A[ramp(i*(int64)2, (int64)1, 2)] = cast(int64, 2 ** 31 - 1) * i;
 }
 The infer result:
 base: int64 -> int64 (since i is involved in another int64 expr)
 stride: int64 -> int32

 Thus ramp should still use int64 for both stride and base after rewrite.
 """
 n = tvm.tir.IntImm("int64", 4)
 m = tvm.tir.IntImm("int64", 2)
 A = te.compute((n, m), lambda i, j: tvm.tir.Cast("int64", 2 ** 31 - 1) * i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[1])
 lower_sch(s, [A], 32, extra_passes=[tvm.tir.transform.VectorizeLoop()])


if __name__ == "__main__":
 test_basic()
 test_thread_axis()
Expand All@@ -263,3 +285,4 @@ def check(shape, index, target_bits, target_dtype):
 test_slice()
 test_relay_basic()
 test_relay_take()
 test_ramp_dtype_consistency()

9 changes: 9 additions & 0 deletions
tests/python/unittest/test_tir_transform_vectorize.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -205,6 +205,14 @@ def test_ir(A, B, C):
 assert expected in error_msg


def test_vectorize_dtype_mismatch():
 n = tvm.tir.IntImm("int64", 4)
 A = te.compute((n,), lambda i: tvm.tir.IntImm("int64", 2 ** 31 - 1) + i, name="A")
 s = te.create_schedule(A.op)
 s[A].vectorize(A.op.axis[0])
 tvm.lower(s, [A], "llvm", simple_mode=True)


if __name__ == "__main__":
 test_vectorize_vector()
 test_vectorize_with_if()
Expand All@@ -214,3 +222,4 @@ def test_ir(A, B, C):
 test_vectorize_with_ge_cond()
 test_vectorize_let()
 test_vectorize_while_fail()
 test_vectorize_dtype_mismatch()