Skip to content

Conversation

@Lunderberg
Copy link
Contributor

This resolves an issue introduced by the combination of #14918 and
#14945. The bug occurred for targets that do not require device-side codegen, but do require a device_type other than kDLCPU. It wasn't caught by CI, as the issue only occurred with the combination of both PRs.

  1. [TIR] SplitHostDevice, handle subroutines #14918 updated SplitHostDevice to only modify the "target" attribute when a device-side function has been extracted.

  2. For VTA, there is no device-side function, as everything is done through host-side API calls.

  3. From (1) and (2), the VTA examples kept the target T.target("ext_dev", host="llvm") after the SplitHostDevice pass, instead of being updated to T.target("llvm").

  4. [TIR] Restrict tir.transform.CombineContextCall to host functions #14945 restricted CombineContextCall to only apply to host-side passes.

  5. From (4) and (5), the CombineContextCall pass was no longer applied to the VTA context calls.

This PR fixes SplitHostDevice, updating the target from T.target("ext_dev", host="llvm") to T.target("llvm"), even if no device sections have been extracted from the function.

This resolves an issue introduced by the combination of
apache#14918 and
apache#14945.  The bug occurred for
targets that do not require device-side codegen, but do require a
`device_type` other than `kDLCPU`.  It wasn't caught by CI, as the
issue only occurred with the combination of both PRs.

1. apache#14918 updated `SplitHostDevice` to only modify the `"target"`
   attribute when a device-side function has been extracted.

2. For VTA, there is no device-side function, as everything is done
   through host-side API calls.

3. From (1) and (2), the VTA examples kept the target
   `T.target("ext_dev", host="llvm")` after the `SplitHostDevice`
   pass, instead of being updated to `T.target("llvm")`.

4. apache#14945 restricted CombineContextCall to only apply to host-side
   passes.

5. From (4) and (5), the `CombineContextCall` pass was no longer
   applied to the VTA context calls.

This PR fixes `SplitHostDevice`, updating the target from
`T.target("ext_dev", host="llvm")` to `T.target("llvm")`, even if no
device sections have been extracted from the function.
@Lunderberg Lunderberg requested a review from tqchen May 30, 2023 14:21
@tvm-bot
Copy link
Collaborator

tvm-bot commented May 30, 2023

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.

Generated by tvm-bot

1 similar comment
@tvm-bot
Copy link
Collaborator

tvm-bot commented May 30, 2023

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.

Generated by tvm-bot

Lunderberg added a commit to Lunderberg/tvm that referenced this pull request May 30, 2023
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
@Lunderberg Lunderberg merged commit 7131411 into apache:main May 30, 2023
@Lunderberg Lunderberg deleted the bugfix_split_host_device branch May 30, 2023 20:43
mei-ye pushed a commit to mei-ye/tvm that referenced this pull request Jun 1, 2023
…pache#14982)

This resolves an issue introduced by the combination of
apache#14918 and
apache#14945.  The bug occurred for
targets that do not require device-side codegen, but do require a
`device_type` other than `kDLCPU`.  It wasn't caught by CI, as the
issue only occurred with the combination of both PRs.

1. apache#14918 updated `SplitHostDevice` to only modify the `"target"`
   attribute when a device-side function has been extracted.

2. For VTA, there is no device-side function, as everything is done
   through host-side API calls.

3. From (1) and (2), the VTA examples kept the target
   `T.target("ext_dev", host="llvm")` after the `SplitHostDevice`
   pass, instead of being updated to `T.target("llvm")`.

4. apache#14945 restricted CombineContextCall to only apply to host-side
   passes.

5. From (4) and (5), the `CombineContextCall` pass was no longer
   applied to the VTA context calls.

This PR fixes `SplitHostDevice`, updating the target from
`T.target("ext_dev", host="llvm")` to `T.target("llvm")`, even if no
device sections have been extracted from the function.
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jun 10, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jun 16, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jun 16, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jun 21, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jul 3, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jul 4, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jul 5, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jul 6, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Jul 7, 2023
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
Lunderberg added a commit to Lunderberg/tvm that referenced this pull request Aug 8, 2023
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
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.

3 participants