[microNPU] enable USMP #10022

manupak · 2022-01-21T19:00:49Z

This commit enables USMP in the microNPU codegen and tests.

manupak · 2022-01-21T19:01:59Z

~~This is blocked on #9929 and #9951~~

* increase test coverage in variants Change-Id: I42a8d7edf3ed92b9643ea01821dbf2f5f97b3e4f

* zephyr cmsis needs more overhead workspace Change-Id: Ifd0071360df4096528d3b1e562e4fa82b1bcb36f

* rebase fixes -- PoolInfo is moved out of usmp namespace Change-Id: Idcf3b4a9a08d8d430eebee9a748ef0f0878379b0

* added more docs Change-Id: I856917523f3d1094d6b5a2335cfd89028f512e37

manupak

I think I've addressed all the comments.

Waiting for CI 🤞

cc: @leandron @Mousius

manupak · 2022-02-17T09:13:57Z

python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py

@@ -105,22 +195,27 @@ def translate(tir_module, params):
    base_addresses : List[util.BaseAddress]
        base addresses to be used by the driver
    """
-
+    candidate_regions_for_scratch = [5, 2, 1]


@ekalda , I did the change as this PR got conflicted. See if it looks good now.

ekalda

LGTM!

ekalda · 2022-02-17T10:33:49Z

python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py

@@ -105,22 +195,27 @@ def translate(tir_module, params):
    base_addresses : List[util.BaseAddress]
        base addresses to be used by the driver
    """
-
+    candidate_regions_for_scratch = [5, 2, 1]


Thanks for adding comment (and the other docstrings as well)! I suppose "regions ranging from 0-6" can be interpreted as both, total of 6 or 7 regions, depending on the counting philosophy :D I think it is clear enough what is meant there though.

Mousius

Slightly worried by the hack I left in being further hacked, thoughts on why that is @manupa-arm ? Otherwise a few naming things that can be looked into later 😸

Mousius · 2022-01-31T09:53:23Z

python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py

+
+    tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+    tvmbaw_region = None


Rather than referencing TVM APIs it's probably better to use workspace_region and workspace_size ?

Hmmm, I was looking for a word differentiate a runtime allocation that is serviced outside of the codegen. Any suggestion?
(workspace_region and workspace_size seems ambigous in that sense)

maybe dynamic_allocation_region ?

Sounds good to me, though low priority 😸

Mousius · 2022-02-17T13:32:37Z

apps/microtvm/zephyr_cmsisnn/src/main.c

@@ -34,7 +34,7 @@ extern float output_storage[12];

 extern const size_t output_len;

-static uint8_t g_crt_workspace[TVMGEN_DEFAULT_WORKSPACE_SIZE + 256];
+static uint8_t g_crt_workspace[TVMGEN_DEFAULT_WORKSPACE_SIZE + 512];


Did I add this? This looks like a hack to increase the workspace size for over-allocation, if we still need this is something broken in USMP?

Yea -- I was puzzled by this too.

In fact, USMP is not enabled for cmsis-nn (yet -- its coming in the next PR) and this is only needed for Zephyr -- other cmsis-nn tests are fine. I am not familiar with Zephyr related impacts here.

i.e. when USMP is enabled we dont need this workspace altogether.

This was surfaced when I removed this :

tvm/python/tvm/micro/model_library_format.py

Lines 217 to 219 in c54a3dd

# TODO(Mousius) - Remove this massive hack when Targets are unified

if target.kind.name in external_codegens:

device_max_workspace[main_target] += int(workspace_size)

.

Which seems to be adding the workspace again for external_codegens.

I think we need to investigate this further

Zephyr should have no real impact on the memory pre-allocated here as it's just a block in flash, this is deeply worrying as the allocator is configured here:

tvm/apps/microtvm/zephyr_cmsisnn/src/main.c

Line 62 in 7d831e1

StackMemoryManager_Init(&app_workspace, g_crt_workspace, TVMGEN_DEFAULT_WORKSPACE_SIZE);

Thus the allocator itself should never go over if it's performing properly, something is very weird here but I agree we should investigate further when we've integrated USMP fully.

Mousius · 2022-02-17T13:35:50Z

python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py

+    (
+        scratch_region_map,
+        tvmbaw_workspace_size,
+        tvmbaw_region,


Similar here, TVMBAW is a detail of how Allocates are lowered

happy to align with the name we end up deciding above...

manupak · 2022-02-17T19:20:57Z

@Mousius ,

As we discussed offline, lets tackle the CMSIS issue in the next PR : #10224 .

For the name suggestion, lets use dynamic_allocation_region and dynamic_allocation_size for which I will do a follow up.

Mousius · 2022-02-21T11:02:59Z

Thanks @manupa-arm 😸 A great step towards unification!

Following an investigation from apache#10022, it turns out, currently the workspace calculation assumes there would be a single lowered PrimFunc could be produced per primitive Relay Function. However, the exception turned out to be the CMSIS-NN codegen that produces multiple calls/PrimFuncs in the place of a single call to single relay PrimFunc. This commit adds changes to workspace calculation to be done on lowered IRModule. Additionally, changes the test utils to not to generate any stack allocator code when USMP is used to make the tests more strict. Change-Id: I5202d9cc7c6a8c00c73791b82df062a8e13dd224

As a follow up to apache#10022, this is a follow PR to perform name change of the region as discussed in that PR. Change-Id: Ifc5244ed2a3b3cc31b089e422c4b3068cd7e26e8

Following an investigation from apache#10022, it turns out, currently the workspace calculation assumes there would be a single lowered PrimFunc could be produced per primitive Relay Function. However, the exception turned out to be the CMSIS-NN codegen that produces multiple calls/PrimFuncs in the place of a single call to single relay PrimFunc. This commit adds changes to workspace calculation to be done on lowered IRModule. Additionally, changes the test utils to not to generate any stack allocator code when USMP is used to make the tests more strict. This change also removes the confusing "run_model" which has semantics identitical to "__tvm_main__" in TIR. Change-Id: I5202d9cc7c6a8c00c73791b82df062a8e13dd224

As a follow up to apache#10022, this is a follow PR to perform name change of the region as discussed in that PR. Change-Id: Ifc5244ed2a3b3cc31b089e422c4b3068cd7e26e8

Following an investigation from apache#10022, it turns out, currently the workspace calculation assumes there would be a single lowered PrimFunc could be produced per primitive Relay Function. However, the exception turned out to be the CMSIS-NN codegen that produces multiple calls/PrimFuncs in the place of a single call to single relay PrimFunc. This commit adds changes to workspace calculation to be done on lowered IRModule. Additionally, changes the test utils to not to generate any stack allocator code when USMP is used to make the tests more strict. This change also removes the confusing "run_model" which has semantics identitical to "__tvm_main__" in TIR. Change-Id: I5202d9cc7c6a8c00c73791b82df062a8e13dd224

Following an investigation from #10022, it turns out, currently the workspace calculation assumes there would be a single lowered PrimFunc could be produced per primitive Relay Function. However, the exception turned out to be the CMSIS-NN codegen that produces multiple calls/PrimFuncs in the place of a single call to single relay PrimFunc. This commit adds changes to workspace calculation to be done on lowered IRModule. Additionally, changes the test utils to not to generate any stack allocator code when USMP is used to make the tests more strict. This change also removes the confusing "run_model" which has semantics identitical to "__tvm_main__" in TIR.

As a follow up to apache#10022, this is a follow PR to perform name change of the region as discussed in that PR. Change-Id: Ifc5244ed2a3b3cc31b089e422c4b3068cd7e26e8

As a follow up to #10022, this is a follow PR to perform name change of the region as discussed in that PR.

This commit enables USMP in the microNPU codegen and tests. The microNPU codegen is modified to support Let nodes that are produced as from USMP.

Following an investigation from apache#10022, it turns out, currently the workspace calculation assumes there would be a single lowered PrimFunc could be produced per primitive Relay Function. However, the exception turned out to be the CMSIS-NN codegen that produces multiple calls/PrimFuncs in the place of a single call to single relay PrimFunc. This commit adds changes to workspace calculation to be done on lowered IRModule. Additionally, changes the test utils to not to generate any stack allocator code when USMP is used to make the tests more strict. This change also removes the confusing "run_model" which has semantics identitical to "__tvm_main__" in TIR.

As a follow up to apache#10022, this is a follow PR to perform name change of the region as discussed in that PR.

manupak marked this pull request as draft January 21, 2022 19:01

manupak force-pushed the micronpu_usmp_v2 branch 2 times, most recently from ff3ca9e to a854d61 Compare January 28, 2022 18:15

manupak marked this pull request as ready for review January 28, 2022 18:16

manupak force-pushed the micronpu_usmp_v2 branch from a854d61 to d709b14 Compare January 28, 2022 22:52

manupak added 4 commits February 17, 2022 07:19

[microNPU] enable USMP

aa31aaa

* increase test coverage in variants Change-Id: I42a8d7edf3ed92b9643ea01821dbf2f5f97b3e4f

[microNPU] enable USMP

76a433c

* zephyr cmsis needs more overhead workspace Change-Id: Ifd0071360df4096528d3b1e562e4fa82b1bcb36f

[microNPU] enable USMP

9daafb1

* rebase fixes -- PoolInfo is moved out of usmp namespace Change-Id: Idcf3b4a9a08d8d430eebee9a748ef0f0878379b0

[microNPU] enable USMP

b7e5d6a

* added more docs Change-Id: I856917523f3d1094d6b5a2335cfd89028f512e37

manupak force-pushed the micronpu_usmp_v2 branch from a33ddf5 to b7e5d6a Compare February 17, 2022 07:45

manupak commented Feb 17, 2022

View reviewed changes

ekalda approved these changes Feb 17, 2022

View reviewed changes

Mousius requested changes Feb 17, 2022

View reviewed changes

Mousius approved these changes Feb 21, 2022

View reviewed changes

Mousius merged commit 2f93780 into apache:main Feb 21, 2022

manupak mentioned this pull request Feb 21, 2022

[AOT] BugFix of workspace calculation #10337

Merged

manupak mentioned this pull request Feb 21, 2022

[microNPU] changing region 'tvmbaw*' to 'dynamic*' #10338

Merged

Mousius pushed a commit that referenced this pull request Mar 18, 2022

[microNPU] changing region 'tvmbaw*' to 'dynamic*' (#10338)

7233c29

As a follow up to #10022, this is a follow PR to perform name change of the region as discussed in that PR.

pfk-beta pushed a commit to pfk-beta/tvm that referenced this pull request Apr 11, 2022

[microNPU] enable USMP (apache#10022)

e56b0f8

This commit enables USMP in the microNPU codegen and tests. The microNPU codegen is modified to support Let nodes that are produced as from USMP.

pfk-beta pushed a commit to pfk-beta/tvm that referenced this pull request Apr 11, 2022

[microNPU] changing region 'tvmbaw*' to 'dynamic*' (apache#10338)

d1dbe67

As a follow up to apache#10022, this is a follow PR to perform name change of the region as discussed in that PR.

driazati mentioned this pull request Jul 14, 2022

TVM v0.9.0.rc0 Release Candidate Notes #12102

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[microNPU] enable USMP #10022

[microNPU] enable USMP #10022

manupak commented Jan 21, 2022

manupak commented Jan 21, 2022 •

edited

Loading

manupak left a comment

manupak Feb 17, 2022

ekalda left a comment

ekalda Feb 17, 2022

Mousius left a comment

Mousius Jan 31, 2022

manupak Feb 17, 2022

manupak Feb 17, 2022

Mousius Feb 21, 2022

Mousius Feb 17, 2022

manupak Feb 17, 2022

manupak Feb 17, 2022

manupak Feb 17, 2022 •

edited

Loading

Mousius Feb 21, 2022

Mousius Feb 17, 2022

manupak Feb 17, 2022

manupak commented Feb 17, 2022 •

edited

Loading

Mousius commented Feb 21, 2022


		tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)

		tvmbaw_region = None

	# TODO(Mousius) - Remove this massive hack when Targets are unified
	if target.kind.name in external_codegens:
	device_max_workspace[main_target] += int(workspace_size)

[microNPU] enable USMP #10022

[microNPU] enable USMP #10022

Conversation

manupak commented Jan 21, 2022

manupak commented Jan 21, 2022 • edited Loading

manupak left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

ekalda left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Mousius left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

manupak Feb 17, 2022 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

manupak commented Feb 17, 2022 • edited Loading

Mousius commented Feb 21, 2022

manupak commented Jan 21, 2022 •

edited

Loading

manupak Feb 17, 2022 •

edited

Loading

manupak commented Feb 17, 2022 •

edited

Loading