elvin-n commented on code in PR #12537:
URL: https://github.com/apache/tvm/pull/12537#discussion_r953584018
##########
python/tvm/topi/adreno/conv2d_nhwc.py:
##########
@@ -270,37 +264,27 @@ def schedule_conv2d_NHWC(cfg, s, output):
##### space definition end #####
pad_data, kernel = s[conv].op.input_tensors
- if (
- isinstance(kernel.op, tvm.te.ComputeOp) and "filter_pack" in
kernel.op.tag
- ): # len(latest.op.axis) == 4:
- # manage scheduling of datacopy
- pad_data, kernel = s[conv].op.input_tensors
+ if autotvm.GLOBAL_SCOPE.in_tuning or input_pack_rt:
if "pad_temp" in pad_data.op.name:
+ s[pad_data].compute_inline()
pack_data = pad_data.op.input_tensors[0]
bind_data_copy(s[pack_data])
else:
bind_data_copy(s[pad_data])
- bind_data_copy(s[kernel])
-
- pad_data, kernel = s[conv].op.input_tensors
- if (
- autotvm.GLOBAL_SCOPE.in_tuning
- or isinstance(kernel.op, tvm.te.ComputeOp)
- and "filter_pack" in kernel.op.tag
- ):
- if "pad_temp" in pad_data.op.name:
- s[pad_data].compute_inline()
AT = s.cache_read(pad_data, get_texture_storage(pad_data.shape),
[conv])
bind_data_copy(s[AT])
- WT = s.cache_read(kernel, get_texture_storage(kernel.shape), [conv])
- bind_data_copy(s[WT])
elif "pad_temp" in pad_data.op.name:
s[pad_data].compute_inline()
# create cache stage
AT = s.cache_read(pad_data, get_texture_storage(pad_data.shape),
[conv])
bind_data_copy(s[AT])
+ if autotvm.GLOBAL_SCOPE.in_tuning or filter_pack_rt:
+ bind_data_copy(s[kernel])
+ WT = s.cache_read(kernel, get_texture_storage(kernel.shape), [conv])
Review Comment:
This is a case of runtime repack of 4d to 5d filter with tir conditions,
i.e.
[here](https://github.com/apache/tvm/blob/main/python/tvm/topi/adreno/utils.py#L243).
Currently it happens if we have output channels not dividable by factor 4. And
this new kernel by will produce buffer. To provide texture to convolution we
are adding cache_read
--
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]
