[microNPU] Integrate the cascader #10862

ekalda · 2022-04-01T12:31:40Z

Integrate the cascader into the codegen and optionally enable it
with the enable_cascader flag. Includes placeholder MemoryRegions until
integration with the PoolInfos provided by a user.

Co-authored-by: Matthew Barrett matthew.barrett@arm.com

ekalda · 2022-04-01T12:44:59Z

cc: @manupa-arm @lhutton1 @NicolaLancellotti

ekalda · 2022-04-01T12:53:19Z

This patch supersedes #10377, so that one can be closed.

manupak

Thanks @ekalda!

It is nice to see this working and green in CI.

I have left few comments around testing and user facing errors.

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

manupak · 2022-04-04T09:37:46Z

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

+
+def _ethos_u55_cascader() -> Callable:
+    flash = MemoryRegion(name="FLASH", size=10 ** 7, read_bandwidth=4, write_bandwidth=4)
+    sram = MemoryRegion(name="SRAM", size=10 ** 6, read_bandwidth=16, write_bandwidth=16)


[Maybe for a subsequent PR] Will it be possible plumb these values using this :

tvm/src/relay/backend/build_module.cc

Lines 414 to 417 in fcdf463

IRModule func_module = WithAttrs(IRModule::FromExpr(func),

{{tvm::attr::kExecutor, executor_},

{tvm::attr::kRuntime, runtime_},

{tvm::attr::kWorkspaceMemoryPools, workspace_memory_pools_}});

An example test :

https://github.com/apache/tvm/blob/main/tests/python/relay/aot/test_crt_aot_usmp.py#L341-L383

Done, with the caveat that we'll assume that there is one workspace pool in the system - is it ok to assume that for now or should we handle the case where there are several workspace pools where some of them are not accessible for the NPU?

manupak · 2022-04-04T09:42:21Z

tests/python/contrib/test_ethosu/test_ops_end_to_end_with_cascader.py

+
+        return conv2d
+
+    infra.compare_tvm_with_tflite(tf_graph, [ifm_shape], accel_type, enable_cascader=True)


For these tests, I think we need to ensure the memory usage is reduced.

We could use calculate_workspace_size TIR analysis utility for that. E.g. :

tvm/tests/python/unittest/test_tir_analysis_calculate_workspace.py

Line 106 in fcdf463

assert tvm.tir.analysis.calculate_workspace_bytes(primfunc, alignment) == size

OR we could use the final memory calculation in a non-USMP flow as follows :

tvm/tests/python/relay/aot/test_crt_aot.py

Lines 937 to 956 in fcdf463

def test_workspace_calculation(workspace_byte_alignment, main_workspace_size):

mod, params = tvm.relay.testing.synthetic.get_workload()

target = "c"

runtime = Runtime("crt")

executor = Executor(

"aot",

{

"workspace-byte-alignment": workspace_byte_alignment,

},

)

with tvm.transform.PassContext(

opt_level=3,

config={

"tir.disable_vectorize": True,

},

):

lib = tvm.relay.build(mod, target, executor=executor, runtime=runtime, params=params)

mlf_memory_map = mlf._build_function_memory_map(lib.function_metadata)

assert mlf_memory_map["main"][0]["workspace_size_bytes"] == main_workspace_size

I changed the tests to check for reduction in memory instead of bitwise accuracy with TFLite (I think it would be still good to have some FVP based tests when cascader is enabled, but it looks like that would need quite a bit of infra refactor, so I'll do it in a separate patch, if this is ok)

lhutton1

Looks good to me modulo @manupa-arm's comments :)

NicolaLancellotti · 2022-04-11T11:35:32Z

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

    return compiler_attrs.accelerator_config


+def enable_cascader():


I think it is better is_cascader_enabled, otherwise, it seems you are enabling it with this function.

Good point! I changed the function name to is_cascader_enabled, but kept the flag/variable as enable_cascader to align with the philosophy of enable_usmp

NicolaLancellotti

LGTM! Thank you @ekalda.

manupak

Broadly looks great!

just a nit and a question

manupak · 2022-04-19T15:16:56Z

tests/python/contrib/test_ethosu/cascader/test_memory_reduction.py

+        mod, params, accel_type, pool_size, enable_cascader=True
+    )
+
+    assert workspace_size_cascader_enabled < workspace_size_cascader_disabled


out of curiosity, should we not check for exact values ?

Yes good point, I changed it to check for the exact values

manupak · 2022-04-19T15:19:33Z

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

+    )
+
+
+def _extract_memory_info(memory_pool):


I think it would be better for this utility to be part of the cascader and not get exposed to the codegen here.

Integrate the cascader into the codegen and optionally enable it with the enable_cascader flag. Includes placeholder MemoryRegions until integration with the PoolInfos provided by a user. Co-authored-by: Matthew Barrett <matthew.barrett@arm.com>

Plumb the workspace memory pools into into the cascader and make the tests to check for the memory reduction.

Change-Id: If2d92846f05a7e8b21be767163841084538805a9

manupak

LGTM with a suggestion for a follow up.

manupak · 2022-04-25T09:06:05Z

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

+        ), "Exactly one workspace pool needs to be provided for the U55 cascader"
+
+        sram = extract_memory_info(workspace_memory_pools.pools[0])
+        tir_mod = LowerToTIR(_ethos_u55_cascader(sram))(mod)


For a followup : Please consider absorbing the call to extract_memory_info inside the cascader. (Sorry for not being clear before). Ideally, we'd want to remove the "MemoryRegion" construct and to get there in the current direction of travel, we should try to confine the usage of it inside the cascader. Therefore the interface of the cascader should be made to accept MemoryPool(s).

I agree that this would be the right way to go about it. Let's do it in a follow up!

manupak · 2022-04-25T09:08:44Z

Thanks @ekalda @mbaret @lhutton1 @NicolaLancellotti !

This is merged now!

* [microNPU] Integrate the cascader Integrate the cascader into the codegen and optionally enable it with the enable_cascader flag. Includes placeholder MemoryRegions until integration with the PoolInfos provided by a user. Co-authored-by: Matthew Barrett <matthew.barrett@arm.com> * Fix linting and a docstring * Plumbing and testing improvements Plumb the workspace memory pools into into the cascader and make the tests to check for the memory reduction. * enable_cascader() -> is_cascader_enabled() * Check for the exact value of workspace size * Remove unused ACCEL_TYPES * Linting... Change-Id: If2d92846f05a7e8b21be767163841084538805a9 * Rebasing... Co-authored-by: Matthew Barrett <matthew.barrett@arm.com>

manupak mentioned this pull request Apr 1, 2022

[microNPU][6] Migrate codegen to use cascader #10377

Closed

manupak requested changes Apr 4, 2022

View reviewed changes

lhutton1 reviewed Apr 4, 2022

View reviewed changes

ekalda force-pushed the integrate-cascader2 branch from 1c20fcd to 2042445 Compare April 11, 2022 11:28

NicolaLancellotti reviewed Apr 11, 2022

View reviewed changes

NicolaLancellotti approved these changes Apr 11, 2022

View reviewed changes

manupak reviewed Apr 19, 2022

View reviewed changes

ekalda force-pushed the integrate-cascader2 branch 4 times, most recently from e764f8c to 094fbcb Compare April 21, 2022 13:46

ekalda and others added 8 commits April 22, 2022 09:30

[microNPU] Integrate the cascader

c3a43b5

Integrate the cascader into the codegen and optionally enable it with the enable_cascader flag. Includes placeholder MemoryRegions until integration with the PoolInfos provided by a user. Co-authored-by: Matthew Barrett <matthew.barrett@arm.com>

Fix linting and a docstring

91ce2ae

Plumbing and testing improvements

c1a0749

Plumb the workspace memory pools into into the cascader and make the tests to check for the memory reduction.

enable_cascader() -> is_cascader_enabled()

55adfe0

Check for the exact value of workspace size

a49cdd0

Remove unused ACCEL_TYPES

31574a7

Linting...

6dcf4b8

Change-Id: If2d92846f05a7e8b21be767163841084538805a9

Rebasing...

f212de6

ekalda force-pushed the integrate-cascader2 branch from 094fbcb to f212de6 Compare April 22, 2022 09:11

manupak approved these changes Apr 25, 2022

View reviewed changes

manupak merged commit d2db9cb into apache:main Apr 25, 2022

manupak mentioned this pull request Jun 30, 2022

[Tracking Issue] Arm(R) Ethos(TM)-U Cascading Scheduler #9429

Closed

12 tasks

driazati mentioned this pull request Jul 14, 2022

TVM v0.9.0.rc0 Release Candidate Notes #12102

Closed

	IRModule func_module = WithAttrs(IRModule::FromExpr(func),
	{{tvm::attr::kExecutor, executor_},
	{tvm::attr::kRuntime, runtime_},
	{tvm::attr::kWorkspaceMemoryPools, workspace_memory_pools_}});


		return conv2d

		infra.compare_tvm_with_tflite(tf_graph, [ifm_shape], accel_type, enable_cascader=True)

	def test_workspace_calculation(workspace_byte_alignment, main_workspace_size):
	mod, params = tvm.relay.testing.synthetic.get_workload()
	target = "c"
	runtime = Runtime("crt")
	executor = Executor(
	"aot",
	{
	"workspace-byte-alignment": workspace_byte_alignment,
	},
	)
	with tvm.transform.PassContext(
	opt_level=3,
	config={
	"tir.disable_vectorize": True,
	},
	):
	lib = tvm.relay.build(mod, target, executor=executor, runtime=runtime, params=params)

	mlf_memory_map = mlf._build_function_memory_map(lib.function_metadata)
	assert mlf_memory_map["main"][0]["workspace_size_bytes"] == main_workspace_size

		return compiler_attrs.accelerator_config


		def enable_cascader():

[microNPU] Integrate the cascader #10862

[microNPU] Integrate the cascader #10862

Uh oh!

Conversation

ekalda commented Apr 1, 2022

Uh oh!

ekalda commented Apr 1, 2022

Uh oh!

ekalda commented Apr 1, 2022

Uh oh!

manupak left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lhutton1 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

NicolaLancellotti left a comment

Choose a reason for hiding this comment

Uh oh!

manupak left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manupak left a comment

Choose a reason for hiding this comment

Uh oh!

manupak Apr 25, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

manupak commented Apr 25, 2022

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

manupak Apr 25, 2022 •

edited

Loading