[HEXAGON] Initial clip operator for Hexagon #11549
Conversation
|
@mehrdadh I'm seeing a segmentation fault in one of the CPU integration steps. Is this a known issue? If so is there a fix? |
|
@TejashShah could you please assign some reviewers to this PR |
|
@jcoplin-quic please add some detail to the PR description. |
|
@mehrdadh I'm not sure why the tests are failing. Would you be able to provide any insight? |
There was a problem hiding this comment.
The error on ci_cpu is because of the missing fixture, but on ci_hexagon it has error on tests/python/contrib/test_hexagon/test_clip.py::TestClipSlice::test_clip_slice[nhwc-8h2w32c2w-2d-input_shape0-output_shape0-nhwc-8h2w32c2w-2d-float16-0.1-0.5]. That is introduced by the this PR.
|
hey @jcoplin-quic, CI was busted when you pushed so the build never started, could you rebase/merge with |
| # build the function | ||
| with tvm.transform.PassContext(opt_level=3): | ||
| func = tvm.build( | ||
| sch, [A, M], tvm.target.Target(target_hexagon, host=target_hexagon), name="clip" |
There was a problem hiding this comment.
The [A, M] argument can be removed. See comment here.
There was a problem hiding this comment.
Hi @Lunderberg removing this argument is causing the build to fail. It's not immediately evident to me why. Would you have any insight?
There was a problem hiding this comment.
Hi @jcoplin-quic, with TIR schedules you'll want to pass in the module, sch.mod, in which case you don't need the tensor list as @Lunderberg pointed out.
There was a problem hiding this comment.
Agreed, the IRModule should be passed to tvm.build, rather than the tvm.tir.Schedule. I've submitted #11913 to improve the error message, as I can definitely see the confusion when the error message doesn't distinguish between expected a te.Schedule or tir.Schedule.
There was a problem hiding this comment.
Thanks @csullivan and @Lunderberg . At line 85, I assign tir_schedule.mod to sch. So sch should be IRModule, not tvm.tir.Schedule
There was a problem hiding this comment.
I think it had to do with optional argument target being treated as a positional argument
jcoplin-quic
force-pushed
the
jcoplin-quic/clip
branch
from
564b1d3 to
f03351f
Compare
* [HEXAGON] Initial clip operator for Hexagon * Changes to utils and infra for pylint * Remove unused import * Use tvm.testing.main() * Address pylint error * Fix incorrect function call * Changes to calls to transform_numpy * Add newline at end of file * Add requires_hexagon and rename under topi * Whitespace fix and reduce input size * Remove te tensor arguments * Correct call to tvm.build * Run black formatting
* [HEXAGON] Initial clip operator for Hexagon * Changes to utils and infra for pylint * Remove unused import * Use tvm.testing.main() * Address pylint error * Fix incorrect function call * Changes to calls to transform_numpy * Add newline at end of file * Add requires_hexagon and rename under topi * Whitespace fix and reduce input size * Remove te tensor arguments * Correct call to tvm.build * Run black formatting
This PR adds a basic Hexagon specific operator and test for clip. Clip is an elementwise operation to restricts an element between two values, min_value and max_value.
This operator is useful for DeepLabv3.
cc @mehrdadh