[Hexagon] Add trivial conv2d operator to Hexagon relay strategy#8915
[Hexagon] Add trivial conv2d operator to Hexagon relay strategy#8915tmoreau89 merged 1 commit intoapache:mainfrom kparzysz-quic:hexagon-trivial-conv2d
Conversation
kparzysz-quic
requested review from
Huyuwei,
Laurawly,
MarisaKirisame,
ZihengJiang,
anijain2305,
jcf94,
jroesch,
junrushao,
jwfromm,
kevinthesun,
masahi,
mbrookhart,
slyubomirsky,
vinx13,
wweic,
yzhliu and
zhiics
as code owners
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]) |
There was a problem hiding this comment.
For Hexagon, don't we need mark pipeline attribute no longer anymore?
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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.
tmoreau89
left a comment
tmoreau89
left a comment
There was a problem hiding this comment.
Thanks @kparzysz-quic ! LGTM
|
Thank you @kparzysz-quic @csullivan - the PR was merged |
…he#8915) This is just to enable relay codegen for conv2d on Hexagon, not for performance.
…he#8915) This is just to enable relay codegen for conv2d on Hexagon, not for performance.
5 participants
This is just to enable relay codegen for conv2d on Hexagon, not for performance.