guberti commented on code in PR #12856:
URL: https://github.com/apache/tvm/pull/12856#discussion_r982171967
##########
python/tvm/topi/arm_cpu/conv2d_alter_op.py:
##########
@@ -121,7 +124,33 @@ def _alter_conv2d_layout(attrs, inputs, tinfos, out_type):
idxd = tvm.tir.indexdiv
- # We don't perform layout alteration for NHWC layout with real data types
+ if topi_tmpl == "depthwise_conv2d_nhwc_dsp.arm_cpu":
+ assert data_layout == "NHWC" and kernel_layout == "HWOI"
+ channels = get_const_tuple(data.shape)[3]
+ KH, KW, _, _ = get_const_tuple(kernel.shape)
+ simd_width = get_dtype_simd_width(data.dtype)
+
+ HWOI_kernel_np = inputs[1].data.numpy()
Review Comment:
Unfortunately, a clean solution is hard, as the strategy function does not
have access to the needed information. When `conv2d_alter_op` is called,
`inputs[1]` (the kernel) has the form:
```python
meta[relay.Constant][0] /* ty=Tensor[(3, 3, 3, 8), int16] */
```
However when the Relay strategy functions are called, `inputs[1]` (the
kernel) looks like:
```python
Tensor(shape=[3, 3, 8, 1], op.name=placeholder)
```
Nowhere inside `relay/op/strategy` do any of the strategy functions check
whether the relevant tensors are constant, so there's not much we can do. I've
added comments explaining this, but please let me know if you have ideas for
how this could be done.
--
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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
