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/06/28 19:40:48 UTC
[GitHub] [tvm] vinx13 opened a new pull request, #11936: [TOPI] Allow conv definition to have custom kernel layout
vinx13 opened a new pull request, #11936:
URL: https://github.com/apache/tvm/pull/11936
This PR added an optional parameter `kernel_layout` to `conv`. This allows arbitrary data and kernel layout combination to be lowered from Relay, and scheduled by auto / meta schedule.
cc @junrushao1994 @masahi @tkonolige
--
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] Ever-Kid commented on a diff in pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
Ever-Kid commented on code in PR #11936:
URL: https://github.com/apache/tvm/pull/11936#discussion_r919593847
##########
python/tvm/topi/nn/conv2d.py:
##########
@@ -57,16 +57,18 @@
)
-def conv2d(input, filter, strides, padding, dilation, layout="NCHW", out_dtype=None):
+def conv2d(
+ input, filter, strides, padding, dilation, data_layout="NCHC", kernel_layout="", out_dtype=None
Review Comment:
maybe data_layout="NCHW" ?
--
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] tkonolige commented on a diff in pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
tkonolige commented on code in PR #11936:
URL: https://github.com/apache/tvm/pull/11936#discussion_r919453757
##########
tests/python/topi/python/test_topi_conv2d_nhwc.py:
##########
@@ -95,5 +95,34 @@ def test_conv2d_nhwc(target, dev, ref_data, dtype, stride, padding, dilation):
tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-5)
+def test_conv2d_explicit_layout(
+ dtype, batch, in_channel, in_size, num_filter, kernel, stride, padding, dilation
+):
+ in_height = in_width = in_size
+ a_shape = (batch, in_height, in_width, in_channel)
+ w_shape = (kernel, kernel, in_channel, num_filter)
+
+ A = te.placeholder(a_shape, name="A", dtype=dtype)
+ W = te.placeholder(w_shape, name="W", dtype=dtype)
+
+ topi.testing.check_compute_definition_equivalent(
+ [
+ A,
+ W,
+ topi.nn.conv2d(
+ A,
+ W,
+ stride,
+ padding,
+ dilation,
+ data_layout="NHWC",
+ kernel_layout="HWIO",
+ out_dtype="float32",
+ ),
+ ],
+ [A, W, topi.nn.conv2d_nhwc(A, W, stride, padding, dilation, "float32")],
Review Comment:
`topi.nn.conv2d_nhwc` just calls `topi.nn.conv2d` under the hood, so this test isn't really checking anything.
--
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] tkonolige commented on pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
tkonolige commented on PR #11936:
URL: https://github.com/apache/tvm/pull/11936#issuecomment-1182533751
@vinx13 I think you can use `tvm.te.create_schedule` as the schedule and the ops will run just fine. Performance might be poor but that doesn't matter for testing correctness.
--
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] vinx13 commented on a diff in pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
vinx13 commented on code in PR #11936:
URL: https://github.com/apache/tvm/pull/11936#discussion_r919455587
##########
tests/python/topi/python/test_topi_conv2d_nhwc.py:
##########
@@ -95,5 +95,34 @@ def test_conv2d_nhwc(target, dev, ref_data, dtype, stride, padding, dilation):
tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-5)
+def test_conv2d_explicit_layout(
+ dtype, batch, in_channel, in_size, num_filter, kernel, stride, padding, dilation
+):
+ in_height = in_width = in_size
+ a_shape = (batch, in_height, in_width, in_channel)
+ w_shape = (kernel, kernel, in_channel, num_filter)
+
+ A = te.placeholder(a_shape, name="A", dtype=dtype)
+ W = te.placeholder(w_shape, name="W", dtype=dtype)
+
+ topi.testing.check_compute_definition_equivalent(
+ [
+ A,
+ W,
+ topi.nn.conv2d(
+ A,
+ W,
+ stride,
+ padding,
+ dilation,
+ data_layout="NHWC",
+ kernel_layout="HWIO",
+ out_dtype="float32",
+ ),
+ ],
+ [A, W, topi.nn.conv2d_nhwc(A, W, stride, padding, dilation, "float32")],
Review Comment:
`topi.nn.conv2d_nhwc` is already covered by existing test cases, but I didn't find any test cases invoke `topi.nn.conv2d` so we just need to check they are the same
Under the hood, both `topi.nn.conv2d_nhwc` and `topi.nn.conv2d` calls `topi.nn.conv` and cover the code path with explicit kernel layout.
--
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] vinx13 commented on pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
vinx13 commented on PR #11936:
URL: https://github.com/apache/tvm/pull/11936#issuecomment-1182503340
@tkonolige There is no good ways to test numerical results for arbitrary combinations of layouts due to lack of schedule. Therefore, I added some tests to check `conv2d` with `data_layout == 'NCHW' and kernel_layotu == 'OIHW'` and `conv2d_nchw` have equivalent compute definition
--
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] Ever-Kid commented on a diff in pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
Ever-Kid commented on code in PR #11936:
URL: https://github.com/apache/tvm/pull/11936#discussion_r919598262
##########
python/tvm/topi/nn/conv2d.py:
##########
@@ -79,17 +81,21 @@ def conv2d(input, filter, strides, padding, dilation, layout="NCHW", out_dtype=N
dilation: int or a list/tuple of two ints
dilation size, or [dilation_height, dilation_width]
- layout : str
+ data_layout : str
layout of data
+ kernel_layout : Optional[str]
+ layout of kernel. If unspecified, use default layout inferred from data_layout. "OHWI" if
+ data_layout == "NCHW", "HWIO" if data_layout == "NHWC".
Review Comment:
A little bit confused to me, shouldn’t "OIHW" & "NCHW" a default pair? I can also see this pattern at line 248.
--
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] Ever-Kid commented on a diff in pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
Ever-Kid commented on code in PR #11936:
URL: https://github.com/apache/tvm/pull/11936#discussion_r919593847
##########
python/tvm/topi/nn/conv2d.py:
##########
@@ -57,16 +57,18 @@
)
-def conv2d(input, filter, strides, padding, dilation, layout="NCHW", out_dtype=None):
+def conv2d(
+ input, filter, strides, padding, dilation, data_layout="NCHC", kernel_layout="", out_dtype=None
Review Comment:
maybe data_layout="NCHW"
--
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] junrushao1994 merged pull request #11936: [TOPI] Allow conv definition to have custom kernel layout
Posted by GitBox <gi...@apache.org>.
junrushao1994 merged PR #11936:
URL: https://github.com/apache/tvm/pull/11936
--
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