[Relay][Topi]Add Sort Op to Relay #6978

mbrookhart · 2020-11-25T21:07:13Z

Re purposing PR for Relay Sort

python/tvm/topi/sort.py

python/tvm/topi/cuda/sort.py

zhiics · 2020-11-30T19:57:31Z

cc @kevinthesun, PTAL when you have time

kevinthesun · 2020-12-01T00:06:13Z

For dynamic topk, can we use Thrust? IMO after sorting with Thrust the only work left is dynamic strided_slice in topi. Not quite sure whether we need to do this in relay level.

anijain2305 · 2020-12-01T00:17:47Z

I tried this PR, I see the following error when used with dynamic input shape

  File "/home/ubuntu/workplace/tvm/t1/tvm/src/relay/backend/compile_engine.cc", line 508
TVMError:
---------------------------------------------------------------
An internal invariant was violated during the execution of TVM.
Please read TVM's error reporting guidelines.
More details can be found here: https://discuss.tvm.ai/t/error-reporting/7793.
---------------------------------------------------------------
  Check failed: fshape_func.count(op) > 0 (0 vs. 0) : Internal error, cannot find ShapeFunc for sort

kevinthesun · 2020-12-01T00:44:35Z

python/tvm/topi/cuda/sort.py

@@ -593,3 +694,82 @@ def schedule_topk(outs):
      The computation schedule for the op.
    """
    return _schedule_sort(outs)
+
+
+def _dyn_topk_legalize(attrs, inputs, arg_types):


I suggest we improve gpu dynamic strided_slice and it helps both dynamic topk and argwhere. For sorting, we can still rely on existing thrust routine.

Any tips for debugging generated code? I have a branch where I've done this a a topi composition, it passes if I compile with AddressSanitizer, but segfaults with a normal build in the generated code. @zhiics attempted to do argwhere, but gets compilation errors with topi dropping variables.

I also get segfault using current topi dynamic strided slice to cut topk result. I'll spend some time on it.

mbrookhart · 2020-12-01T17:49:17Z

@anijain2305 I'll add a shape func for sort and a regression test

@kevinthesun I added a thrust implementation of sort, and tested dynamic topk with thrust enabled, it worked just fine.

kevinthesun · 2020-12-03T00:07:57Z

#7018
Fix for dynamic strided slice in topi to enable topk.

mbrookhart · 2020-12-07T19:54:42Z

closing in favor of #7018

mbrookhart · 2020-12-08T20:50:01Z

@masahi @zhiics @jwfromm

@kevinthesun found a fix for dynamic topk, so I don't think I need this legalization anymore, but do you think the sort op at the relay level is still useful?

masahi · 2020-12-08T21:00:25Z

yes hummingbird people have asked for a converter for pytorch sort op, that would be exactly what I need. Thanks.

cc @interesaaat

mbrookhart · 2020-12-08T21:01:28Z

Okay, I'll reopen and rip out the legalization.

mbrookhart · 2020-12-09T19:40:05Z

@kevinthesun Can you re-review without the topk legalization?

python/tvm/topi/cuda/sort.py

masahi · 2020-12-16T22:27:08Z

python/tvm/topi/cuda/sort.py

+        lambda ins, outs: tvm.tir.call_packed(
+            "tvm.contrib.thrust.sort", ins[0], outs[0], outs[1], is_ascend
+        ),
+        out_buffers=[value_buf, indices_buf],


There is no need to pass indices_buf for normal sort, but tvm.contrib.thrust.sort always expects indices_buf to be passed in... What we call tvm.contrib.thrust.sort is really argsort.

For optimal performance we should have the true tvm.contrib.thrust.sort someday.

Hmm, you're right, we don't strictly need that. Want me to add a todo?

Yeah that would be nice

The review is stale

masahi

Please fix remaining doc issues, otherwise LGTM.

masahi · 2020-12-17T22:00:45Z

Thanks @mbrookhart @zhiics @kevinthesun @jwfromm

* Add sort op to relay * fix lint * fix sort docstring * fix docs * add TODO, shape_func, cleanup * add dynamic tests for sort and argsort

apivovarov · 2022-06-11T21:52:35Z

python/tvm/topi/cuda/sort.py

+    Parameters
+    ----------
+    outs: Array of Tensor
+        The computation graph description of argsort


typo: argsort -> sort

jwfromm reviewed Nov 25, 2020

View reviewed changes

python/tvm/topi/sort.py Outdated Show resolved Hide resolved

zhiics reviewed Nov 25, 2020

View reviewed changes

python/tvm/topi/cuda/sort.py Outdated Show resolved Hide resolved

zhiics reviewed Nov 30, 2020

View reviewed changes

python/tvm/topi/cuda/sort.py Outdated Show resolved Hide resolved

zhiics approved these changes Nov 30, 2020

View reviewed changes

kevinthesun previously requested changes Dec 1, 2020

View reviewed changes

mbrookhart closed this Dec 7, 2020

mbrookhart deleted the sort branch December 7, 2020 19:54

mbrookhart restored the sort branch December 8, 2020 21:01

mbrookhart reopened this Dec 8, 2020

mbrookhart force-pushed the sort branch from 5943a65 to 49f0a48 Compare December 8, 2020 21:07

mbrookhart changed the title ~~[Relay][Topi][Dynamic] Add a Sort op and use the legalization pass to perform dynamic topk on GPU~~ [Relay][Topi]Add Sort Op to Relay Dec 8, 2020

masahi reviewed Dec 16, 2020

View reviewed changes

python/tvm/topi/cuda/sort.py Outdated Show resolved Hide resolved

Matthew Brookhart added 3 commits December 16, 2020 12:13

Add sort op to relay

73e49c1

fix lint

b20e671

fix sort docstring

d67cfdd

mbrookhart force-pushed the sort branch from b1ee86d to d67cfdd Compare December 16, 2020 19:13

masahi reviewed Dec 16, 2020

View reviewed changes

python/tvm/topi/cuda/sort.py Outdated Show resolved Hide resolved

masahi reviewed Dec 16, 2020

View reviewed changes

python/tvm/topi/cuda/sort.py Outdated Show resolved Hide resolved

masahi reviewed Dec 16, 2020

View reviewed changes

masahi approved these changes Dec 16, 2020

View reviewed changes

mbrookhart added 3 commits December 16, 2020 15:39

fix docs

ea16f08

add TODO, shape_func, cleanup

79cee17

add dynamic tests for sort and argsort

2202b84

masahi merged commit 9e766d9 into apache:main Dec 17, 2020

masahi pushed a commit to masahi/tvm that referenced this pull request Dec 18, 2020

[Relay][Topi]Add Sort Op to Relay (apache#6978)

91a5aa1

* Add sort op to relay * fix lint * fix sort docstring * fix docs * add TODO, shape_func, cleanup * add dynamic tests for sort and argsort

trevor-m pushed a commit to neo-ai/tvm that referenced this pull request Jan 21, 2021

[Relay][Topi]Add Sort Op to Relay (apache#6978)

c3a6951

* Add sort op to relay * fix lint * fix sort docstring * fix docs * add TODO, shape_func, cleanup * add dynamic tests for sort and argsort

junrushao mentioned this pull request Nov 1, 2021

Apache TVM v0.8 Release Note Candidate #9416

Closed

apivovarov reviewed Jun 11, 2022

View reviewed changes

apivovarov added a commit to neo-ai/tvm that referenced this pull request Jun 12, 2022

Add register_shape_func for argsort (apache#6978)

8227062

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Relay][Topi]Add Sort Op to Relay #6978

[Relay][Topi]Add Sort Op to Relay #6978

mbrookhart commented Nov 25, 2020 •

edited

zhiics commented Nov 30, 2020

kevinthesun commented Dec 1, 2020 •

edited

anijain2305 commented Dec 1, 2020

kevinthesun Dec 1, 2020

mbrookhart Dec 1, 2020 •

edited

kevinthesun Dec 1, 2020

mbrookhart commented Dec 1, 2020

kevinthesun commented Dec 3, 2020

mbrookhart commented Dec 7, 2020

mbrookhart commented Dec 8, 2020

masahi commented Dec 8, 2020

mbrookhart commented Dec 8, 2020

mbrookhart commented Dec 9, 2020

masahi Dec 16, 2020

mbrookhart Dec 16, 2020

masahi Dec 16, 2020

masahi left a comment

masahi commented Dec 17, 2020

apivovarov Jun 11, 2022

[Relay][Topi]Add Sort Op to Relay #6978

[Relay][Topi]Add Sort Op to Relay #6978

Conversation

mbrookhart commented Nov 25, 2020 • edited

zhiics commented Nov 30, 2020

kevinthesun commented Dec 1, 2020 • edited

anijain2305 commented Dec 1, 2020

kevinthesun Dec 1, 2020

Choose a reason for hiding this comment

mbrookhart Dec 1, 2020 • edited

Choose a reason for hiding this comment

kevinthesun Dec 1, 2020

Choose a reason for hiding this comment

mbrookhart commented Dec 1, 2020

kevinthesun commented Dec 3, 2020

mbrookhart commented Dec 7, 2020

mbrookhart commented Dec 8, 2020

masahi commented Dec 8, 2020

mbrookhart commented Dec 8, 2020

mbrookhart commented Dec 9, 2020

masahi Dec 16, 2020

Choose a reason for hiding this comment

mbrookhart Dec 16, 2020

Choose a reason for hiding this comment

masahi Dec 16, 2020

Choose a reason for hiding this comment

masahi left a comment

Choose a reason for hiding this comment

masahi commented Dec 17, 2020

apivovarov Jun 11, 2022

Choose a reason for hiding this comment

mbrookhart commented Nov 25, 2020 •

edited

kevinthesun commented Dec 1, 2020 •

edited

mbrookhart Dec 1, 2020 •

edited