[TIR] Restrict tir.transform.InstallDebugSpans to host functions #14943

Lunderberg · 2023-05-24T20:35:24Z

Previously, the tir.InstallDebugSpans pass required the module to contain only a single PrimFunc. This commit relaxes the requirement, to require a single host-side PrimFunc, and to ignore any other device-side functions.

tvm-bot · 2023-05-24T20:35:27Z

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.

cc @Hzfengsy, @junrushao, @quic-sanirudh, @shingjan _{See #10317 for details}

_{Generated by tvm-bot}

Lunderberg · 2023-05-24T20:41:39Z

PRs #14943, #14944, and #14945 perform analogous changes to three different host-only lowering passes. Each of these three contains the commits that add utility functions tvm::tir::IsHostFunc and TargetNode::HasKey, then apply the new utility to a specific pass. Whichever of the three lands first can have the others rebased on top to prevent any conflicts.

These PRs in turn are a subset of the functionality in #14862. As that PR had become unmanageably large, I am carving it out into several smaller independent PRs for ease of review.

junrushao

tbh i dont know if this pass is actually used but lets get it in

Lunderberg · 2023-05-25T12:36:24Z

tbh i dont know if this pass is actually used but lets get it in

It looks like it requires the "tir.enable_debug" option in the PassConfig. I'm not sure how often that gets enabled, but it was added relatively recently (~5 months ago by @driazati ), so I'm guessing there's some usage of it.

There's also some potential improvements, since an entire IRModule can be printed out, which could probably remove the host-only and one-function-only limitations, but I didn't want to mix functionality changes with implementation changes at the moment.

Lunderberg · 2023-05-25T12:36:30Z

@tvm-bot rerun

PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict.

This utility method makes it easier to determine if a target contains a specific key.

For modules that contain both host and device functions, this utility function checks whether a given PrimFunc is a host function, based on the target annotation.

Previously, the `tir.InstallDebugSpans` pass required the module to contain only a single PrimFunc. This commit relaxes the requirement, to require a single host-side PrimFunc, and to ignore any other device-side functions.

Lunderberg · 2023-05-25T14:46:37Z

Rebased onto #14950 to allow CI to run while waiting for the bugfix to land.

The functionality tested in this commit was added across several recent PRs, each of which tested their features in isolation. This PR adds unit tests to validate the end-to-end behavior of TIR subroutine calls. PRs building up to this point: - TVMScript - apache#14889 - apache#14915 - apache#14919 - apache#14941 - Functionality improvements of existing TIR passes - apache#14913 - apache#14914 - apache#14918 - apache#14951 - Changes to the TIR lowering flow - apache#14942 - apache#14985 - Codegen updates - apache#14958 - apache#14901 - Compatibility updates/fixes - apache#14892 - apache#14950 - apache#14943 - apache#14944 - apache#14945 - apache#14952 - apache#14982 - apache#14949

…che#14943) * [Bugfix] Avoid symbol conflicts in MakePackedAPI/MakeUnpackedAPI PRs apache#14913 and apache#14914 made analogous changes to `MakePackedAPI` and `MakeUnpackedAPI` to handle subroutine calls. Both PRs introduced the same symbol, `tvm::tir::SubroutineCallRewriter`, a local utility to update internal calls to a modified function. While each PR passed CI individually, and was therefore able to merge, having both changes caused a duplicate symbol. This commit updates `MakePackedAPI` and `MakeUnpackedAPI` to place their local utilities into anonymous namespaces, avoiding the conflict. * [Target] Added utility method TargetNode::HasKey() This utility method makes it easier to determine if a target contains a specific key. * [TIR] Added utility method tvm::tir::IsHostFunc(const PrimFunc&) For modules that contain both host and device functions, this utility function checks whether a given PrimFunc is a host function, based on the target annotation. * [TIR] Restrict InstallDebugSpans to host functions Previously, the `tir.InstallDebugSpans` pass required the module to contain only a single PrimFunc. This commit relaxes the requirement, to require a single host-side PrimFunc, and to ignore any other device-side functions.

The functionality tested in this commit was added across several recent PRs, each of which tested their features in isolation. This PR adds unit tests to validate the end-to-end behavior of TIR subroutine calls. PRs building up to this point: - TVMScript - apache#14889 - apache#14915 - apache#14919 - apache#14941 - Functionality improvements of existing TIR passes - apache#14913 - apache#14914 - apache#14918 - apache#14951 - Changes to the TIR lowering flow - apache#14942 - apache#14985 - Codegen updates - apache#14958 - apache#14901 - Compatibility updates/fixes - apache#14892 - apache#14950 - apache#14943 - apache#14944 - apache#14945 - apache#14952 - apache#14982 - apache#14949

This was referenced May 24, 2023

[TIR] Restrict tir.transform.LowerTVMBuiltin to host functions #14944

Merged

[TIR] Restrict tir.transform.CombineContextCall to host functions #14945

Merged

junrushao approved these changes May 24, 2023

View reviewed changes

Lunderberg added 4 commits May 25, 2023 09:36

[Target] Added utility method TargetNode::HasKey()

8c27815

This utility method makes it easier to determine if a target contains a specific key.

[TIR] Added utility method tvm::tir::IsHostFunc(const PrimFunc&)

0918064

For modules that contain both host and device functions, this utility function checks whether a given PrimFunc is a host function, based on the target annotation.

[TIR] Restrict InstallDebugSpans to host functions

7b47f29

Previously, the `tir.InstallDebugSpans` pass required the module to contain only a single PrimFunc. This commit relaxes the requirement, to require a single host-side PrimFunc, and to ignore any other device-side functions.

Lunderberg force-pushed the restrict_installdebugspan_to_host branch from 1a762db to 7b47f29 Compare May 25, 2023 14:46

masahi merged commit 875217c into apache:main May 26, 2023

Lunderberg deleted the restrict_installdebugspan_to_host branch May 26, 2023 11:30

ysh329 mentioned this pull request Jul 12, 2023

[Release] v0.13.0 Release Candidate Notes #15295

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[TIR] Restrict tir.transform.InstallDebugSpans to host functions #14943

[TIR] Restrict tir.transform.InstallDebugSpans to host functions #14943

Lunderberg commented May 24, 2023

tvm-bot commented May 24, 2023

Lunderberg commented May 24, 2023

junrushao left a comment

Lunderberg commented May 25, 2023

Lunderberg commented May 25, 2023

Lunderberg commented May 25, 2023

[TIR] Restrict tir.transform.InstallDebugSpans to host functions #14943

[TIR] Restrict tir.transform.InstallDebugSpans to host functions #14943

Conversation

Lunderberg commented May 24, 2023

tvm-bot commented May 24, 2023

Lunderberg commented May 24, 2023

junrushao left a comment

Choose a reason for hiding this comment

Lunderberg commented May 25, 2023

Lunderberg commented May 25, 2023

Lunderberg commented May 25, 2023