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

[Hexagon] Add trivial conv2d operator to Hexagon relay strategy #8915

Merged
merged 1 commit into from
Sep 3, 2021
Merged

[Hexagon] Add trivial conv2d operator to Hexagon relay strategy #8915

merged 1 commit into from
Sep 3, 2021

Conversation

kparzysz-quic
Copy link
Contributor

This is just to enable relay codegen for conv2d on Hexagon, not for performance.

This is just to enable relay codegen for conv2d on Hexagon, not for
performance.

def schedule_conv2d_nhwc(outs):
"""Schedule for Conv2d NHWC operator."""
s = tvm.te.create_schedule([x.op for x in outs])
Copy link
Member

Choose a reason for hiding this comment

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

For Hexagon, don't we need mark pipeline attribute no longer anymore?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

We're moving towards running entire graphs on Hexagon, so no offloading is necessary. You can still write an op that will use pipeline, but for relay-based compilation we will use ops that run entirely on Hexagon.

Copy link
Contributor

Choose a reason for hiding this comment

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

Can we add a possibility to appoint a layer to be offloaded to CPU? I have compared TFLite case and SNPE one for object detection models and it happened that it's better to execute Detection Output on CPU instead of DSP. It might be an exception but such use case is quite widely used

Copy link
Contributor Author

@kparzysz-quic kparzysz-quic Sep 3, 2021

Choose a reason for hiding this comment

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

Qualcomm ships products that do this sort of a thing, but I'm not familiar with how the separation between CPU and DSP happens there. In these products, the communication between CPU and DSP is handled by the runtime that is a part of these products, and the offloading does not occur within an individual op. That means that even there there is no need to offload anything explicitly in TVM, because the entire op will run either on CPU or on DSP.

Edit: I think I misunderstood your question---yes SNPE does allow that, but we are not planning to implement this in TVM in the whole-graph compilation mode, at least not in the near term.

Copy link
Contributor

Choose a reason for hiding this comment

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

I might have misunderstood your message We're moving towards running entire graphs on Hexagon, so no offloading is necessary. My point that sometimes it might be more efficient to offload some parts and give control to user and I got impression that we are going to cancel this ability.

Copy link
Contributor Author

@kparzysz-quic kparzysz-quic Sep 3, 2021

Choose a reason for hiding this comment

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

Yes, but we're not planning to support this in the near term.

Edit: we don't have that ability in TVM (for Hexagon) right now, so we're not canceling anything.

Copy link
Contributor

@csullivan csullivan left a comment

Choose a reason for hiding this comment

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

LGTM

Copy link
Contributor

@tmoreau89 tmoreau89 left a comment

Choose a reason for hiding this comment

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

Thanks @kparzysz-quic ! LGTM

@tmoreau89 tmoreau89 merged commit 318e3fe into apache:main Sep 3, 2021
@tmoreau89
Copy link
Contributor

Thank you @kparzysz-quic @csullivan - the PR was merged

@kparzysz-quic kparzysz-quic deleted the hexagon-trivial-conv2d branch September 4, 2021 12:02
ylc pushed a commit to ylc/tvm that referenced this pull request Sep 29, 2021
…he#8915)

This is just to enable relay codegen for conv2d on Hexagon, not for
performance.
ylc pushed a commit to ylc/tvm that referenced this pull request Jan 13, 2022
…he#8915)

This is just to enable relay codegen for conv2d on Hexagon, not for
performance.
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.

5 participants