Skip to content
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.

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

[Bug] Crash when applying transform_layout with axis separator #13974

Closed
quic-sanirudh opened this issue Feb 13, 2023 · 2 comments
Closed

[Bug] Crash when applying transform_layout with axis separator #13974

quic-sanirudh opened this issue Feb 13, 2023 · 2 comments
Labels
needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug

Comments

@quic-sanirudh
Copy link
Contributor

When we apply transform_layout on a write buffer generated by cache_write with axis separators, we see an error that the sref pointed by the block does not exist in stmt2ref. The below code reproduces the error:

Expected behavior

The code should work fine without errors and apply the layout transform with axis separators

Actual behavior

Code crashes with error Check failed: (it != self_->stmt2ref.end()) is false: InternalError: The sref points to a statement that does not exist in stmt2ref

Environment

Tested on Upstream TVM main with Ubuntu 18.04

Steps to reproduce

import tvm
import tvm.contrib.hexagon
import tvm.testing
from tvm.script import tir as T


@T.prim_func
def add(
    p0: T.Buffer((T.int64(33), T.int64(128)), "float32"),
    p1: T.Buffer((T.int64(33), T.int64(128)), "float32"),
    T_add: T.Buffer((T.int64(33), T.int64(128)), "float32"),
):
    T.func_attr({"global_symbol": "main", "tir.noalias": True})
    # with T.block("root"):
    for ax0, ax1 in T.grid(T.int64(33), T.int64(128)):
        with T.block("T_add"):
            v_ax0, v_ax1 = T.axis.remap("SS", [ax0, ax1])
            T.reads(p0[v_ax0, v_ax1], p1[v_ax0, v_ax1])
            T.writes(T_add[v_ax0, v_ax1])
            T_add[v_ax0, v_ax1] = p0[v_ax0, v_ax1] + p1[v_ax0, v_ax1]


def transform_fn(x, y):
    return [x // 32, y, tvm.te.AXIS_SEPARATOR, x % 32]


def main():
    sch = tvm.tir.Schedule(add, debug_mask="all")
    block_rv = sch.get_block("T_add")
    sch.cache_write(block_rv, 0, "global")
    sch.transform_layout(block_rv, ("write", 0), transform_fn, pad_value=0.0)


if __name__ == "__main__":
    main()

Triage

  • needs-triage
@quic-sanirudh quic-sanirudh added needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug labels Feb 13, 2023
@quic-sanirudh
Copy link
Contributor Author

cc @Lunderberg @vinx13

quic-sanirudh added a commit to quic-sanirudh/tvm that referenced this issue Feb 17, 2023
A mismatch between the blocks present in the `result` vs the blocks
passed in block_sref_to_reuse caused the bug mentioned in apache#13974. This
patch tries to fix that bug by collecting only the blocks that are part
of result and also present in the block replacement map
new_block_to_old_
quic-sanirudh added a commit to quic-sanirudh/tvm that referenced this issue Feb 17, 2023
A mismatch between the blocks present in the `result` vs the blocks
passed in block_sref_to_reuse caused the bug mentioned in apache#13974. This
patch tries to fix that bug by collecting only the blocks that are part
of result and also present in the block replacement map
new_block_to_old_
quic-sanirudh added a commit to quic-sanirudh/tvm that referenced this issue Feb 17, 2023
A mismatch between the blocks present in the `result` vs the blocks
passed in `block_sref_to_reuse` caused the bug mentioned in apache#13974.

This patch tries to fix that bug by collecting only the blocks that are
part of result and also present in the block replacement map
`new_block_to_old_`. Since the scope block is `result`, only that block
and its child blocks would be replaced, and any replaced block would be
present in `rewriter.new_block_to_old_`. Thus, collecting the replaced
blocks from among child blocks of `result` guarantees that the
`block_sref_reuse` would contain all the replaced blocks and that
they'll point to the correct block in `result` thus avoiding the missing
SRef error.
tqchen pushed a commit that referenced this issue Feb 18, 2023
* [TIR] [Bugfix] Pass the correct block_sref_reuse to Replace

A mismatch between the blocks present in the `result` vs the blocks
passed in `block_sref_to_reuse` caused the bug mentioned in #13974.

This patch tries to fix that bug by collecting only the blocks that are
part of result and also present in the block replacement map
`new_block_to_old_`. Since the scope block is `result`, only that block
and its child blocks would be replaced, and any replaced block would be
present in `rewriter.new_block_to_old_`. Thus, collecting the replaced
blocks from among child blocks of `result` guarantees that the
`block_sref_reuse` would contain all the replaced blocks and that
they'll point to the correct block in `result` thus avoiding the missing
SRef error.
@quic-sanirudh
Copy link
Contributor Author

The issue is fixed with #14023, hence closing this issue.

yongwww pushed a commit to yongwww/tvm that referenced this issue Feb 27, 2023
…4023)

* [TIR] [Bugfix] Pass the correct block_sref_reuse to Replace

A mismatch between the blocks present in the `result` vs the blocks
passed in `block_sref_to_reuse` caused the bug mentioned in apache#13974.

This patch tries to fix that bug by collecting only the blocks that are
part of result and also present in the block replacement map
`new_block_to_old_`. Since the scope block is `result`, only that block
and its child blocks would be replaced, and any replaced block would be
present in `rewriter.new_block_to_old_`. Thus, collecting the replaced
blocks from among child blocks of `result` guarantees that the
`block_sref_reuse` would contain all the replaced blocks and that
they'll point to the correct block in `result` thus avoiding the missing
SRef error.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
needs-triage PRs or issues that need to be investigated by maintainers to find the right assignees to address it type: bug
Projects
None yet
Development

No branches or pull requests

1 participant