Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[TOPI] Allow conv definition to have custom kernel layout #11936

Merged
merged 4 commits into from
Jul 14, 2022

Conversation

vinx13
Copy link
Member

@vinx13 vinx13 commented Jun 28, 2022

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

Copy link
Member

@junrushao junrushao left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

Copy link
Contributor

@tkonolige tkonolige left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks pretty good, but I don't see any tests for different kernel_layouts. Could you add some?

@vinx13
Copy link
Member Author

vinx13 commented Jul 12, 2022

@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

out_dtype="float32",
),
],
[A, W, topi.nn.conv2d_nhwc(A, W, stride, padding, dilation, "float32")],
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

topi.nn.conv2d_nhwc just calls topi.nn.conv2d under the hood, so this test isn't really checking anything.

Copy link
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

@tkonolige
Copy link
Contributor

@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.

@vinx13 vinx13 force-pushed the feat/conv_generic_kernel_layout branch 3 times, most recently from c8daa85 to 1f1b428 Compare July 13, 2022 00:28
@@ -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
Copy link
Contributor

@Ever-Kid Ever-Kid Jul 13, 2022

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

maybe data_layout="NCHW" ?

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".
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

A little bit confused to me, shouldn’t "OIHW" & "NCHW" a default pair? I can also see this pattern at line 248.

Copy link
Contributor

@tkonolige tkonolige left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good! Thanks for adding the tests!

@vinx13 vinx13 force-pushed the feat/conv_generic_kernel_layout branch from cf01ebb to 249be49 Compare July 13, 2022 20:40
@vinx13 vinx13 force-pushed the feat/conv_generic_kernel_layout branch from 249be49 to a9eb3d4 Compare July 13, 2022 21:15
@junrushao junrushao merged commit a571bfb into apache:main Jul 14, 2022
xinetzone pushed a commit to daobook/tvm that referenced this pull request Nov 25, 2022
* [TOPI] Allow conv definition to have custom kernel layout

* add tests

* fix

* fix
mikeseven pushed a commit to mikeseven/tvm that referenced this pull request Sep 27, 2023
* [TOPI] Allow conv definition to have custom kernel layout

* add tests

* fix

* fix
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants