junrushao commented on code in PR #13195:
URL: https://github.com/apache/tvm/pull/13195#discussion_r1016005557
##########
python/tvm/topi/nn/conv2d.py:
##########
@@ -989,6 +989,119 @@ def unpack_NCHWc_to_nchw(packed_out, out_dtype):
return unpacked_out
[email protected]_func
+def conv2d_winograd_nhwc(
+ data,
+ weight,
+ strides,
+ padding,
+ dilation,
+ out_dtype,
+ pre_computed=False,
+ auto_scheduler_rewritten_layout="",
+ meta_schedule_original_shape=None,
+):
+ """Conv2D Winograd in NHWC layout.
+ This is a clean version to be used by the auto-scheduler for both CPU and
GPU.
+
+ Parameters
+ ----------
+ data : tvm.te.Tensor
+ 4-D with shape [batch, in_height, in_width, in_channel]
+ weight : tvm.te.Tensor
+ 4-D with shape [filter_height, filter_width, in_channel, num_filter]
+ strides : int or a list/tuple of two ints
+ stride size, or [stride_height, stride_width]
+ padding : int or a list/tuple of two ints
+ padding size, or [pad_height, pad_width]
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+ out_dtype : str, optional
+ Specifies the output data type.
+ pre_computed: bool
+ Whether the kernel is precomputed
+ auto_scheduler_rewritten_layout: str = ""
+ The layout after auto-scheduler's layout rewrite pass.
+ meta_schedule_original_shape: Optional[List[PrimExpr]] = None
+ The original shape of the input tensor.
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 4-D with shape [batch, out_height, out_width, out_channel]
+ """
+ tile_size = 4
+ return _conv2d_winograd_nhwc_impl(
+ data,
+ weight,
+ strides,
+ padding,
+ dilation,
+ out_dtype,
+ tile_size,
+ pre_computed=pre_computed,
+ write_cache_level=2,
+ auto_scheduler_rewritten_layout=auto_scheduler_rewritten_layout,
+ meta_schedule_original_shape=meta_schedule_original_shape,
+ )
+
+
[email protected]_func
+def conv2d_winograd_nchw(
+ data,
+ weight,
+ strides,
+ padding,
+ dilation,
+ out_dtype,
+ pre_computed=False,
+ auto_scheduler_rewritten_layout="",
+ meta_schedule_original_shape=None,
+):
+ """Conv2D Winograd in NCHW layout.
+ This is a clean version to be used by the auto-scheduler for both CPU and
GPU.
+
+ Parameters
+ ----------
+ data : tvm.te.Tensor
+ 4-D with shape [batch, in_channel, in_height, in_width]
+ weight : tvm.te.Tensor
+ 4-D with shape [filter_height, filter_width, in_channel, num_filter]
+ strides : int or a list/tuple of two ints
+ stride size, or [stride_height, stride_width]
+ padding : int or a list/tuple of two ints
+ padding size, or [pad_height, pad_width]
+ dilation: int or a list/tuple of two ints
+ dilation size, or [dilation_height, dilation_width]
+ out_dtype : str, optional
+ Specifies the output data type.
+ pre_computed: bool
+ Whether the kernel is precomputed
+ auto_scheduler_rewritten_layout: str = ""
+ The layout after auto-scheduler's layout rewrite pass.
+ meta_schedule_original_shape: Optional[List[PrimExpr]] = None
+ The original shape of the input tensor.
+
+ Returns
+ -------
+ output : tvm.te.Tensor
+ 4-D with shape [batch, out_height, out_width, out_channel]
+ """
+ tile_size = 4
Review Comment:
not for simplicity but for consistency with auto-scheduler's nhwc case a
couple of lines above. note that this configuration is also overridden by cuda
dispatch, so we dont have to worry too much about 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: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
