You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/09/30 19:58:30 UTC
[GitHub] [tvm] adstraw opened a new pull request, #12954: [Hexagon] 3-stage pipeline; multi queue async DMA for cache read / write
adstraw opened a new pull request, #12954:
URL: https://github.com/apache/tvm/pull/12954
Add `HexagonUserDMA` support for multiple virtual queues which enables Async DMA for both `cache_read` and `cache_write` while maintaining a single descriptor chain to maintain overall FIFO ordering between virtual queues. Tested with runtime unit tests and at the python level for a simple operator.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] adstraw commented on pull request #12954: [Hexagon] 3-stage pipeline; multi queue async DMA for cache read / write
Posted by GitBox <gi...@apache.org>.
adstraw commented on PR #12954:
URL: https://github.com/apache/tvm/pull/12954#issuecomment-1263973994
CC @masahi
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] adstraw commented on a diff in pull request #12954: [Hexagon] 3-stage pipeline; multi queue async DMA for cache read / write
Posted by GitBox <gi...@apache.org>.
adstraw commented on code in PR #12954:
URL: https://github.com/apache/tvm/pull/12954#discussion_r986049998
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -46,18 +47,33 @@ def plus_one_ref(a):
return plus_one_primfunc, plus_one_ref
-@tvm.testing.requires_hexagon
-def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inner, dtype, scope):
+@tvm.testing.fixture
Review Comment:
This is an artifact of the GitHub diff. I split out the schedule as a parameter here. The test now starts on line 74.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] adstraw commented on a diff in pull request #12954: [Hexagon] 3-stage pipeline; multi queue async DMA for cache read / write
Posted by GitBox <gi...@apache.org>.
adstraw commented on code in PR #12954:
URL: https://github.com/apache/tvm/pull/12954#discussion_r986077448
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -46,18 +47,33 @@ def plus_one_ref(a):
return plus_one_primfunc, plus_one_ref
-@tvm.testing.requires_hexagon
-def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inner, dtype, scope):
+@tvm.testing.fixture
+def schedule(compute, sched, scope):
sch = tir.Schedule(compute[0])
- root = sch.get_block("root")
+
compute_block = sch.get_block("compute")
cache_read_block = sch.cache_read(compute_block, 0, scope)
i, _ = sch.get_loops(compute_block)
sch.compute_at(cache_read_block, i)
- sch.annotate(i, "software_pipeline_stage", [0, 1])
- sch.annotate(i, "software_pipeline_order", [0, 1])
- sch.annotate(i, "software_pipeline_async_stages", [0])
+
+ if sched == "cache_read":
Review Comment:
Good idea. I can add this test.
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -46,18 +47,33 @@ def plus_one_ref(a):
return plus_one_primfunc, plus_one_ref
-@tvm.testing.requires_hexagon
-def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inner, dtype, scope):
+@tvm.testing.fixture
+def schedule(compute, sched, scope):
sch = tir.Schedule(compute[0])
- root = sch.get_block("root")
+
compute_block = sch.get_block("compute")
cache_read_block = sch.cache_read(compute_block, 0, scope)
i, _ = sch.get_loops(compute_block)
sch.compute_at(cache_read_block, i)
- sch.annotate(i, "software_pipeline_stage", [0, 1])
- sch.annotate(i, "software_pipeline_order", [0, 1])
- sch.annotate(i, "software_pipeline_async_stages", [0])
+
+ if sched == "cache_read":
Review Comment:
Good idea. I can add this test. Can either do it here or in a future PR.
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -26,8 +26,9 @@
outer = tvm.testing.parameter(8, 16)
inner = tvm.testing.parameter(64, 128)
-scope = tvm.testing.parameter("global", "global.vtcm")
dtype = tvm.testing.parameter("uint8", "float16")
+scope = tvm.testing.parameter("global", "global.vtcm")
+sched = tvm.testing.parameter("cache_read", "cache_read_write")
@tvm.testing.fixture
Review Comment:
Correct. This allows for any number of `cache_read` and `cache_write` stages to be lowered to Async DMA. Note that there is a known issue when trying to do `cache_read` for an op with multiple inputs in the same stage which will be addressed in a future PR. Future PR will modify compute on this test to be `a + b` instead of `a + 1` and add support to use Async DMA to `cache_read` both `a` and `b`.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [tvm] tmoreau89 commented on a diff in pull request #12954: [Hexagon] 3-stage pipeline; multi queue async DMA for cache read / write
Posted by GitBox <gi...@apache.org>.
tmoreau89 commented on code in PR #12954:
URL: https://github.com/apache/tvm/pull/12954#discussion_r986037292
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -46,18 +47,33 @@ def plus_one_ref(a):
return plus_one_primfunc, plus_one_ref
-@tvm.testing.requires_hexagon
-def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inner, dtype, scope):
+@tvm.testing.fixture
+def schedule(compute, sched, scope):
sch = tir.Schedule(compute[0])
- root = sch.get_block("root")
+
compute_block = sch.get_block("compute")
cache_read_block = sch.cache_read(compute_block, 0, scope)
i, _ = sch.get_loops(compute_block)
sch.compute_at(cache_read_block, i)
- sch.annotate(i, "software_pipeline_stage", [0, 1])
- sch.annotate(i, "software_pipeline_order", [0, 1])
- sch.annotate(i, "software_pipeline_async_stages", [0])
+
+ if sched == "cache_read":
Review Comment:
Does it make sense for the sake of completeness to test a cache_write pipeline where `sch.annotate(i, "software_pipeline_async_stages", [1])`?
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -46,18 +47,33 @@ def plus_one_ref(a):
return plus_one_primfunc, plus_one_ref
-@tvm.testing.requires_hexagon
-def test_software_pipeline_with_cache_read(hexagon_launcher, compute, outer, inner, dtype, scope):
+@tvm.testing.fixture
Review Comment:
is `requires_hexagon` now implicit?
##########
tests/python/contrib/test_hexagon/test_software_pipeline_async.py:
##########
@@ -26,8 +26,9 @@
outer = tvm.testing.parameter(8, 16)
inner = tvm.testing.parameter(64, 128)
-scope = tvm.testing.parameter("global", "global.vtcm")
dtype = tvm.testing.parameter("uint8", "float16")
+scope = tvm.testing.parameter("global", "global.vtcm")
+sched = tvm.testing.parameter("cache_read", "cache_read_write")
@tvm.testing.fixture
Review Comment:
With this PR, we could technically test any n-stage pipeline, correct (doesn't have to be limited to 3-stage)?
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org