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/01/21 19:01:01 UTC
[GitHub] [tvm] manupa-arm opened a new pull request #10022: [microNPU] enable USMP
manupa-arm opened a new pull request #10022:
URL: https://github.com/apache/tvm/pull/10022
This commit enables USMP in the microNPU codegen and tests.
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1043328116
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.
--
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] Hzfengsy commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1030531864
Hi @manupa-arm. Thanks for the explanation. That makes sense to me
--
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] Mousius commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Mousius commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r811011385
##########
File path: 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];
Review comment:
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:
https://github.com/apache/tvm/blob/7d831e1dd34d441e9217096d0baf78e6f97c77d9/apps/microtvm/zephyr_cmsisnn/src/main.c#L62
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.
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809061775
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
Review comment:
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)
--
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] ekalda commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806999093
##########
File path: 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]
Review comment:
Thanks for the clarification, makes sense! Yeah, it makes sense to do it in the follow-up since the CI is green...
--
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] Hzfengsy commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Hzfengsy commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1030515810
I don't know about the codebase, but shall we add a switch in CMake files to let users decide whether enable it?
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r799405283
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
+ if len(candidate_regions_for_scratch) > 0:
+ tvmbaw_region = candidate_regions_for_scratch.pop()
+
+ # Need a mutable data structure to be updated by the following function
+ # Therefore, using a list instead of int
+ tvmbaw_size = [0]
Review comment:
Cool! sounds like a good idea
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r808825500
##########
File path: 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]
Review comment:
@ekalda , I did the change as this PR got conflicted. See if it looks good now.
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809059560
##########
File path: 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];
Review comment:
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.
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1030521761
@Hzfengsy
I don't think so.
We generally do that if that component has a external dependency. USMP is developed natively developed as a TIR pass in TVM. This commit just uses in our codegen which has a CMake config anyway.
I find that It is similar to asking should we add a CMake variable to enable TIR vectorizer just because some flows dont use it.
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809066670
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -106,21 +202,31 @@ def translate(tir_module, params):
base addresses to be used by the driver
"""
+ # The NPU has 6 usable regions ranging from 0-6
+ # The regions 0, 3, and 4 is already used for input,
+ # output and constant, respectively (See _get_regions()).
+ # Thus, for scratch we are left with 5, 2 and 1.
+ candidate_regions_for_scratch = [5, 2, 1]
+ (
+ scratch_region_map,
+ tvmbaw_workspace_size,
+ tvmbaw_region,
Review comment:
happy to align with the name we end up deciding above...
--
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] Mousius commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Mousius commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r811013071
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
Review comment:
Sounds good to me, though low priority :smile_cat:
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1018780014
This is blocked on #9929 and #9951
--
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] lhutton1 commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
lhutton1 commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r797451191
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
+ if len(candidate_regions_for_scratch) > 0:
+ tvmbaw_region = candidate_regions_for_scratch.pop()
+
+ # Need a mutable data structure to be updated by the following function
+ # Therefore, using a list instead of int
+ tvmbaw_size = [0]
Review comment:
It might be clearer to use the `nonlocal` keyword to keep this value an int e.g.
```
j = 0
def something():
nonlocal j
j += 1
```
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1040025471
A friendly ping to get this moving..
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806982398
##########
File path: 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]
Review comment:
Good point! I ll add a comment, in a follow up -- if thats alright ?
The comment would be something to describe the following point :
We have regions from 0 to 5 and we already are using regions 0, 3, and 4 statically for constants, input and output, respectively. Therefore, it leaves us with three more regions : 5, 2 and 1 to be used by memory pools.
##########
File path: 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]
Review comment:
Good point! I ll add a comment, in a follow up -- if thats alright ?
The comment would be something to describe the following point :
We have regions from 0 to 5 and we already are using regions 0, 3, and 4 statically for constants, input and output, respectively (see static_regions above). Therefore, it leaves us with three more regions : 5, 2 and 1 to be used by memory pools.
--
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] manupa-arm edited a comment on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm edited a comment on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1043328116
@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.
--
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] ekalda commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806892717
##########
File path: 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]
Review comment:
What is the meaning of these numbers? Maybe a comment would help there...
##########
File path: 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]
Review comment:
Thanks for the clarification, makes sense! Yeah, it makes sense to do it in the follow-up since the CI is green...
--
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] ekalda commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806892717
##########
File path: 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]
Review comment:
What is the meaning of these numbers? Maybe a comment would help there...
--
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] Mousius commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Mousius commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809051534
##########
File path: 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];
Review comment:
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?
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -106,21 +202,31 @@ def translate(tir_module, params):
base addresses to be used by the driver
"""
+ # The NPU has 6 usable regions ranging from 0-6
+ # The regions 0, 3, and 4 is already used for input,
+ # output and constant, respectively (See _get_regions()).
+ # Thus, for scratch we are left with 5, 2 and 1.
+ candidate_regions_for_scratch = [5, 2, 1]
+ (
+ scratch_region_map,
+ tvmbaw_workspace_size,
+ tvmbaw_region,
Review comment:
Similar here, TVMBAW is a detail of how Allocates are lowered
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
Review comment:
Rather than referencing TVM APIs it's probably better to use `workspace_region` and `workspace_size` ?
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809064416
##########
File path: 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];
Review comment:
This was surfaced when I removed this : https://github.com/apache/tvm/blob/c54a3dd6999ead14fd118e6d84f2b64a1ecf3b9d/python/tvm/micro/model_library_format.py#L217-L219.
Which seems to be adding the workspace again for external_codegens.
I think we need investigate this further
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809064416
##########
File path: 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];
Review comment:
This was surfaced when I removed this : https://github.com/apache/tvm/blob/c54a3dd6999ead14fd118e6d84f2b64a1ecf3b9d/python/tvm/micro/model_library_format.py#L217-L219.
Which seems to be adding the workspace again for external_codegens.
I think we need to investigate this further
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1025943879
cc : @ekalda @lhutton1 @mbaret @Mousius for reviews.
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1040025471
A friendly ping to get this moving..
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809060176
##########
File path: 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];
Review comment:
i.e. when USMP is enabled we dont need this workspace altogether.
--
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] manupa-arm edited a comment on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm edited a comment on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1018780014
~~This is blocked on #9929 and #9951~~
--
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] manupa-arm commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1029924692
a friendly ping for reviews @mbaret @Mousius @ekalda
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806982398
##########
File path: 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]
Review comment:
Good point! I ll add a comment, in a follow up -- if thats alright ?
The comment would be something to describe the following point :
We have regions from 0 to 5 and we already are using regions 0, 3, and 4 statically for constants, input and output, respectively (see static_regions above). Therefore, it leaves us with three more regions : 5, 2 and 1 to be used by memory pools.
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r806982398
##########
File path: 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]
Review comment:
Good point! I ll add a comment, in a follow up -- if thats alright ?
The comment would be something to describe the following point :
We have regions from 0 to 5 and we already are using regions 0, 3, and 4 statically for constants, input and output, respectively. Therefore, it leaves us with three more regions : 5, 2 and 1 to be used by memory pools.
--
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] manupa-arm commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
manupa-arm commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r809067781
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -81,6 +72,107 @@ def get_accelerator_arch_config(accel_type):
return accel_config_str_map[accel_type]
+class RegionOffset(NamedTuple):
+ """A data structure to hold region and address offset corresponding to a tensor"""
+
+ region: int
+ offset: int
+
+
+def analyze_scratch_memory_acesses(mod: tvm.IRModule, candidate_regions_for_scratch: List[int]):
+ """
+ Parameters
+ ----------
+ mod: tvm.IRModule
+ The TIR module containing ethosu extern calls
+ candidate_regions_for_scratch: List[int]
+ A list of region integers that could be used for scratch regions
+
+ Returns
+ -------
+ scratch_region_map : Dict[tvm.tir.Var, int]
+ A map between buffer vars to scratch regions they are assigned
+ tvm_backend_alloc_workspace_size : int
+ The size of tvm_backend_alloc_workspace call required to service
+ remaining allocate nodes if any
+ tvm_backend_alloc_workspace_region : int
+ The region associated with the tvm_backend_alloc_workspace
+ """
+ scratch_region_map = dict()
+ pool_var_region_map = dict()
+ # There should only be a single function
+ assert len(mod.functions.items()) == 1
+ primfunc = mod.functions.items()[0][1]
+ if "pool_args" in primfunc.attrs.keys():
+ pool_args = primfunc.attrs["pool_args"]
+ for pool_arg in pool_args:
+ pool_param = primfunc.params[int(pool_arg.pool_var_idx)]
+ pool_var_region_map[pool_param] = candidate_regions_for_scratch.pop()
+ scratch_region_map[pool_param] = RegionOffset(
+ region=pool_var_region_map[pool_param], offset=None
+ )
+
+ def analyze_pool_access(stmt):
+ if isinstance(stmt, tvm.tir.stmt.LetStmt):
+ call_address_of = stmt.value
+ load = call_address_of.args[0]
+ pool_var = load.buffer_var
+ scratch_region_map[stmt.var] = RegionOffset(
+ region=pool_var_region_map[pool_var], offset=int(load.index)
+ )
+
+ tvm.tir.stmt_functor.post_order_visit(primfunc.body, analyze_pool_access)
+
+ tvmbaw_region = None
Review comment:
maybe dynamic_allocation_region ?
--
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] ekalda commented on a change in pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
ekalda commented on a change in pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#discussion_r808902750
##########
File path: 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]
Review comment:
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.
--
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] Mousius merged pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Mousius merged pull request #10022:
URL: https://github.com/apache/tvm/pull/10022
--
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] Mousius commented on pull request #10022: [microNPU] enable USMP
Posted by GitBox <gi...@apache.org>.
Mousius commented on pull request #10022:
URL: https://github.com/apache/tvm/pull/10022#issuecomment-1046752647
Thanks @manupa-arm :smile_cat: A great step towards unification!
--
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