Skip to content

Conversation

@wrongtest-intellif
Copy link
Contributor

Fix the following script with multiple kernel launch (each take standalone thread binding) which is not round-trip-able.

import tvm
from tvm.script import tir as T

@T.prim_func
def main(A: T.Buffer[128, "float32"], C: T.Buffer[128, "float32"]):
    B = T.alloc_buffer([128], "float32")
    for i in range(128):
        with T.block("B1"):
            vi = T.axis.remap("S", [i])
            B[vi] = A[vi] + 1.0
    for i in range(128):
        with T.block("B2"):
            vi = T.axis.remap("S", [i])
            C[vi] = B[vi] + 2.0

s = tvm.tir.schedule.Schedule(main)
b1 = s.get_block("B1")
b2 = s.get_block("B2")
i = s.get_loops(b1)[0]
s.bind(i, "threadIdx.x")
i = s.get_loops(b2)[0]
s.bind(i, "threadIdx.x")
f = tvm.lower(s.mod["main"])
print(f.script())
tvm.ir.assert_structural_equal(f, tvm.script.from_source(f.script()))

Since the env thread is always printed on the function head, TryDeallocVar which decrease the name uniqueness counting may not be invoked for thread binding code path.

@tvm-bot
Copy link
Collaborator

tvm-bot commented Dec 15, 2022

Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from Reviewers by @-ing them in a comment.

  • No users to tag found in teams: tvmscript See #10317 for details

Generated by tvm-bot

Copy link
Member

@Hzfengsy Hzfengsy left a comment

Choose a reason for hiding this comment

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

LGTM. Thanks for the fixing

@vinx13 vinx13 merged commit ce97138 into apache:main Dec 15, 2022
fzi-peccia pushed a commit to fzi-peccia/tvm that referenced this pull request Mar 27, 2023
…#13622)

* Fix print round-tripable multi thread env binding

* add unittest
mikeseven pushed a commit to mikeseven/tvm that referenced this pull request Sep 27, 2023
…#13622)

* Fix print round-tripable multi thread env binding

* add unittest
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants