Skip to content

[TIR][Schedule] Allow named block and buffer arguments in Schedule#11624

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
vinx13 merged 3 commits into from
Jun 9, 2022
Merged

[TIR][Schedule] Allow named block and buffer arguments in Schedule #11624

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
c6a43f9
[Schedule] Allowed string argument as block arg
Lunderberg Jun 8, 2022
07594c2
[Schedule] Allowed string argument to Schedule.reindex
Lunderberg Jun 8, 2022
209139e
CI test update
Lunderberg Jun 8, 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
112 changes: 76 additions & 36 deletions python/tvm/tir/schedule/schedule.py
View file
Open in desktop

Large diffs are not rendered by default.

9 changes: 4 additions & 5 deletions src/tir/schedule/primitive/cache_read_write.cc
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
Integer buffer_index_type) {
PythonAPICall py("reindex");
py.Input("block", block);
py.Input("buffer_index", buffer_index);
py.Input("buffer_index_type", '"' +
std::string(BufferIndexType2Str(
static_cast<BufferIndexType>(buffer_index_type->value))) +
'"');
std::ostringstream os;
os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
<< "\", " << buffer_index << ")";
py.Input("buffer", os.str());
py.SingleOutput(outputs);
return py.Str();
}
Expand Down
94 changes: 48 additions & 46 deletions tests/python/unittest/test_tir_schedule_cache_read_write.py
View file
Open in desktop
Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_c = sch.get_block("C")
cached_a = sch.cache_read(block_b, 0, "global")
cached_b = sch.cache_read(block_c, 0, "local")
cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
sch = tir.Schedule(access_under_scope, debug_mask="all")
block_b = sch.get_block("B")
block_c = sch.get_block("C")
block_b = "B" if use_block_name else sch.get_block("B")
block_c = "C" if use_block_name else sch.get_block("C")
sch.cache_read(block_b, 0, "local")
sch.cache_read(block_c, 0, "global")
tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
sch = tir.Schedule(opaque_access, debug_mask="all")
block = sch.get_block("load_store")
block = "load_store" if use_block_name else sch.get_block("load_store")
sch.cache_read(block, 0, "global")
tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
sch = tir.Schedule(func_multi_consumer, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
sch.cache_read(block_b, 0, "global")
tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_c = sch.get_block("C")
block_c = "C" if use_block_name else sch.get_block("C")
sch.cache_read(block_c, 0, "shared")
sch.cache_read(block_c, 0, "local")
tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
block = sch.get_block("consumer")
block = "consumer" if use_block_name else sch.get_block("consumer")
sch.cache_read(block, 0, "shared")
tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
sch.cache_read(block_b, 0, "global")
tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
sch = tir.Schedule(func_multi_producer, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_c = sch.get_block("C")
cached_b = sch.cache_write(block_b, 0, "local")
cached_c = sch.cache_write(block_c, 0, "global")
cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
sch = tir.Schedule(access_under_scope, debug_mask="all")
block_a = sch.get_block("A")
block_b = sch.get_block("B")
block_a = "A" if use_block_name else sch.get_block("A")
block_b = "B" if use_block_name else sch.get_block("B")
block_scope = sch.get_block("scope")
sch.cache_write(block_a, 0, "local")
sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
sch = tir.Schedule(opaque_access, debug_mask="all")
block_store = sch.get_block("load_store")
block_opaque = sch.get_block("opaque")
block_match_buffer = sch.get_block("match_buffer")
block_store = "load_store" if use_block_name else sch.get_block("load_store")
block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
sch.cache_write(block_store, 0, "global")
sch.cache_write(block_opaque, 0, "global")
sch.cache_write(block_match_buffer, 0, "global")
tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
sch = tir.Schedule(func_multi_consumer, debug_mask="all")
block_a = sch.get_block("A")
block_a = "A" if use_block_name else sch.get_block("A")
sch.cache_write(block_a, 0, "global")
tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
sch.cache_write(block_b, 0, "shared")
sch.cache_write(block_b, 0, "local")
tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
# cache write for intermediate buffer
sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
block = sch.get_block("producer")
block = "producer" if use_block_name else sch.get_block("producer")
sch.cache_write(block, 0, "shared")
tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
# cache write for external buffer
sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
block = sch.get_block("consumer")
block = "consumer" if use_block_name else sch.get_block("consumer")
sch.cache_write(block, 0, "shared")
tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
sch = tir.Schedule(func_multi_producer, debug_mask="all")
block_a0 = sch.get_block("A0")
block_a1 = sch.get_block("A1")
block_a0 = "A0" if use_block_name else sch.get_block("A0")
block_a1 = "A1" if use_block_name else sch.get_block("A1")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_write(block_a0, 0, "global")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
sch = tir.Schedule(elementwise, debug_mask="all")
block_b = sch.get_block("B")
block_b = "B" if use_block_name else sch.get_block("B")
with pytest.raises(tvm.tir.ScheduleError):
sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading
, 'i'); if (__m === '*' || __re.test(location.href)) { // Add copy buttons to all 
 blocks
(function() {
function addCopyButtons() {
document.querySelectorAll('pre code').forEach(function(codeBlock) {
if (codeBlock.parentElement.hasAttribute('data-copy-added')) return;
codeBlock.parentElement.setAttribute('data-copy-added', 'true');
var btn = document.createElement('button');
btn.textContent = 'Copy';
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;';
btn.onmouseover = function() { this.style.opacity = '1'; };
btn.onmouseout = function() { this.style.opacity = '0.7'; };
btn.onclick = function() {
navigator.clipboard.writeText(codeBlock.textContent).then(function() {
btn.textContent = 'Copied!';
setTimeout(function() { btn.textContent = 'Copy'; }, 1500);
});
};
codeBlock.parentElement.style.position = 'relative';
codeBlock.parentElement.appendChild(btn);
});
}
addCopyButtons();
// Re-run on dynamic content
var observer = new MutationObserver(addCopyButtons);
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:Add Copy Buttons to Code Blocks]', __e); }
})();
(function(){
try {
var __m = "github.com";
var __re = new RegExp('^' + "github\\.com" + '
[TIR][Schedule] Allow named block and buffer arguments in Schedule by Lunderberg · Pull Request #11624 · apache/tvm · GitHub
Skip to content
[TIR][Schedule] Allow named block and buffer arguments in Schedule#11624

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Force GitHub README to respect dark mode
(function() {
var style = document.createElement('style');
style.textContent = '
.markdown-body {
color-scheme: dark light;
}
.markdown-body pre { background: #161b22 !important; }
.markdown-body code { background: rgba(110, 118, 129, 0.4) !important; }
.markdown-body table th, .markdown-body table td { border-color: #30363d !important; }
.markdown-body img { background: #0d1117; }
.markdown-body blockquote { border-left-color: #8b949e; }
.markdown-body hr { border-color: #30363d; }
';
document.head.appendChild(style);
})();
}
} catch(__e) { console.warn('[Userscript:GitHub Dark Mode README Fix]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
[TIR][Schedule] Allow named block and buffer arguments in Schedule by Lunderberg · Pull Request #11624 · apache/tvm · GitHub
Skip to content
[TIR][Schedule] Allow named block and buffer arguments in Schedule#11624

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

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

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

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

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Auto-enable theater mode on YouTube
(function() {
function tryTheater() {
var btn = document.querySelector('button[aria-label="Theater mode"], ytd-player #player button[title="Theater mode"]');
if (btn && !btn.classList.contains('activated')) {
btn.click();
}
}
// Try immediately
tryTheater();
// Try after navigation (SPA)
var lastUrl = location.href;
setInterval(function() {
if (location.href !== lastUrl) {
lastUrl = location.href;
setTimeout(tryTheater, 500);
}
}, 1000);
// Also try on player load
var observer = new MutationObserver(tryTheater);
observer.observe(document.body, { childList: true, subtree: true });
})();
}
} catch(__e) { console.warn('[Userscript:YouTube Theater Mode Default]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
[TIR][Schedule] Allow named block and buffer arguments in Schedule by Lunderberg · Pull Request #11624 · apache/tvm · GitHub
Skip to content
[TIR][Schedule] Allow named block and buffer arguments in Schedule#11624

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

, 'i');
if (__m === '*' || __re.test(location.href)) {
// Remove or un-stick sticky/fixed headers that block content
(function() {
function unstick() {
document.querySelectorAll('header, nav, [role="banner"], .header, .navbar, .sticky, .fixed-top, [style*="position: fixed"], [style*="position:sticky"]').forEach(function(el) {
if (el.style.position === 'fixed' || el.style.position === 'sticky' || getComputedStyle(el).position === 'fixed' || getComputedStyle(el).position === 'sticky') {
el.style.position = 'static';
el.style.top = 'auto';
el.style.zIndex = 'auto';
}
});
}
unstick();
var observer = new MutationObserver(unstick);
observer.observe(document.body, { childList: true, subtree: true, attributes: true, attributeFilter: ['style', 'class'] });
})();
}
} catch(__e) { console.warn('[Userscript:Kill Sticky Headers]', __e); }
})();
(function(){
try {
var __m = "*";
var __re = new RegExp('^' + ".*" + '
[TIR][Schedule] Allow named block and buffer arguments in Schedule by Lunderberg · Pull Request #11624 · apache/tvm · GitHub
Skip to content
[TIR][Schedule] Allow named block and buffer arguments in Schedule#11624

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading

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

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

vinx13
merged 3 commits into

from

Jun 9, 2022
 Merged


[TIR][Schedule] Allow named block and buffer arguments in Schedule
#11624
Show file tree
Hide file tree
 Changes from all commits
Commits

Show all changes


3 commits


Select commit
Hold shift + click to select a range
c6a43f9

[Schedule] Allowed string argument as block arg


Lunderberg Jun 8, 2022
07594c2

[Schedule] Allowed string argument to Schedule.reindex


Lunderberg Jun 8, 2022
209139e

CI test update


Lunderberg Jun 8, 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


112 changes: 76 additions & 36 deletions
python/tvm/tir/schedule/schedule.py

View file


Open in desktop


Large diffs are not rendered by default.


9 changes: 4 additions & 5 deletions
src/tir/schedule/primitive/cache_read_write.cc

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -1241,11 +1241,10 @@ struct ReIndexTraits : public UnpackedInstTraits<ReIndexTraits> {
 Integer buffer_index_type) {
 PythonAPICall py("reindex");
 py.Input("block", block);
 py.Input("buffer_index", buffer_index);
 py.Input("buffer_index_type", '"' +
 std::string(BufferIndexType2Str(
 static_cast<BufferIndexType>(buffer_index_type->value))) +
 '"');
 std::ostringstream os;
 os << "(\"" << BufferIndexType2Str(static_cast<BufferIndexType>(buffer_index_type->value))
 << "\", " << buffer_index << ")";
 py.Input("buffer", os.str());
 py.SingleOutput(outputs);
 return py.Str();
 }
Expand Down

94 changes: 48 additions & 46 deletions
tests/python/unittest/test_tir_schedule_cache_read_write.py

View file


Open in desktop

Original file line numberDiff line numberDiff line change
Expand Up@@ -741,13 +741,15 @@ def block_predicate_cache_write_output_buf() -> None:

########## Testcases for cache_read ##########

use_block_name = tvm.testing.parameter(by_dict={"block_obj": False, "block_name": True})

def test_cache_read_elementwise():

def test_cache_read_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_a = sch.cache_read(block_b, 0, "global")
 cached_b = sch.cache_read(block_c, 0, "local")
 cached_a = sch.cache_read("B" if use_block_name else block_b, 0, "global")
 cached_b = sch.cache_read("C" if use_block_name else block_c, 0, "local")
 assert sch.get(cached_a) == sch.get(sch.get_block("A_global"))
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -756,87 +758,87 @@ def test_cache_read_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_under_scope():
def test_cache_read_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_b, 0, "local")
 sch.cache_read(block_c, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_under_scope, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_read_opaque_access():
def test_cache_read_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block = sch.get_block("load_store")
 block = "load_store" if use_block_name else sch.get_block("load_store")
 sch.cache_read(block, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_read_location():
def test_cache_read_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_read():
def test_continuous_cache_read(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_c = sch.get_block("C")
 block_c = "C" if use_block_name else sch.get_block("C")
 sch.cache_read(block_c, 0, "shared")
 sch.cache_read(block_c, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_read_with_block_predicate():
def test_cache_read_with_block_predicate(use_block_name):
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_read(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_read, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_read_non_int32_shape():
def test_cache_read_non_int32_shape(use_block_name):
 sch = tir.Schedule(elementwise_shape_int64, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_read(block_b, 0, "global")
 tvm.ir.assert_structural_equal(cache_read_shape_int64, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise_shape_int64)


def test_cache_read_fail_multi_producer():
def test_cache_read_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "global")


def test_cache_read_fail_index_out_of_bound():
def test_cache_read_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 1, "global")


def test_cache_read_fail_invalid_storage_scope():
def test_cache_read_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_read(block_b, 0, "test_scope")


########## Testcases for cache_write ##########


def test_cache_write_elementwise():
def test_cache_write_elementwise(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_c = sch.get_block("C")
 cached_b = sch.cache_write(block_b, 0, "local")
 cached_c = sch.cache_write(block_c, 0, "global")
 cached_b = sch.cache_write("B" if use_block_name else block_b, 0, "local")
 cached_c = sch.cache_write("C" if use_block_name else block_c, 0, "global")
 assert sch.get(cached_b) == sch.get(sch.get_block("B_local"))
 assert sch.get(cached_c) == sch.get(sch.get_block("C_global"))
 assert sch.get(block_b) == sch.get(sch.get_block("B"))
Expand All@@ -845,10 +847,10 @@ def test_cache_write_elementwise():
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_under_scope():
def test_cache_write_under_scope(use_block_name):
 sch = tir.Schedule(access_under_scope, debug_mask="all")
 block_a = sch.get_block("A")
 block_b = sch.get_block("B")
 block_a = "A" if use_block_name else sch.get_block("A")
 block_b = "B" if use_block_name else sch.get_block("B")
 block_scope = sch.get_block("scope")
 sch.cache_write(block_a, 0, "local")
 sch.cache_write(block_b, 0, "global")
Expand All@@ -857,70 +859,70 @@ def test_cache_write_under_scope():
 verify_trace_roundtrip(sch=sch, mod=access_under_scope)


def test_cache_write_opaque_access():
def test_cache_write_opaque_access(use_block_name):
 sch = tir.Schedule(opaque_access, debug_mask="all")
 block_store = sch.get_block("load_store")
 block_opaque = sch.get_block("opaque")
 block_match_buffer = sch.get_block("match_buffer")
 block_store = "load_store" if use_block_name else sch.get_block("load_store")
 block_opaque = "opaque" if use_block_name else sch.get_block("opaque")
 block_match_buffer = "match_buffer" if use_block_name else sch.get_block("match_buffer")
 sch.cache_write(block_store, 0, "global")
 sch.cache_write(block_opaque, 0, "global")
 sch.cache_write(block_match_buffer, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_opaque_access, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=opaque_access)


def test_cache_write_location():
def test_cache_write_location(use_block_name):
 sch = tir.Schedule(func_multi_consumer, debug_mask="all")
 block_a = sch.get_block("A")
 block_a = "A" if use_block_name else sch.get_block("A")
 sch.cache_write(block_a, 0, "global")
 tvm.ir.assert_structural_equal(cache_write_multi_consumer, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_multi_consumer)


def test_continuous_cache_write():
def test_continuous_cache_write(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 sch.cache_write(block_b, 0, "shared")
 sch.cache_write(block_b, 0, "local")
 tvm.ir.assert_structural_equal(continuous_cache_write, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=elementwise)


def test_cache_write_with_block_predicate():
def test_cache_write_with_block_predicate(use_block_name):
 # cache write for intermediate buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("producer")
 block = "producer" if use_block_name else sch.get_block("producer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_intermediate_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)
 # cache write for external buffer
 sch = tir.Schedule(func_with_block_predicate, debug_mask="all")
 block = sch.get_block("consumer")
 block = "consumer" if use_block_name else sch.get_block("consumer")
 sch.cache_write(block, 0, "shared")
 tvm.ir.assert_structural_equal(block_predicate_cache_write_output_buf, sch.mod["main"])
 verify_trace_roundtrip(sch=sch, mod=func_with_block_predicate)


def test_cache_write_fail_multi_producer():
def test_cache_write_fail_multi_producer(use_block_name):
 sch = tir.Schedule(func_multi_producer, debug_mask="all")
 block_a0 = sch.get_block("A0")
 block_a1 = sch.get_block("A1")
 block_a0 = "A0" if use_block_name else sch.get_block("A0")
 block_a1 = "A1" if use_block_name else sch.get_block("A1")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a0, 0, "global")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_a1, 0, "global")


def test_cache_write_fail_index_out_of_bound():
def test_cache_write_fail_index_out_of_bound(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 1, "global")


def test_cache_write_fail_invalid_storage_scope():
def test_cache_write_fail_invalid_storage_scope(use_block_name):
 sch = tir.Schedule(elementwise, debug_mask="all")
 block_b = sch.get_block("B")
 block_b = "B" if use_block_name else sch.get_block("B")
 with pytest.raises(tvm.tir.ScheduleError):
 sch.cache_write(block_b, 0, "test_scope")

Expand Down
Loading