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/01/24 15:07:23 UTC
[GitHub] [tvm] jinhongyii opened a new pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
jinhongyii opened a new pull request #10043:
URL: https://github.com/apache/tvm/pull/10043
This PR is one of the schedule rule for MetaSchedule.
The rule does not support auto tensorization for now.
Co-authored-by: Junru Shao junrushao1994@gmail.com
Co-authored-by: Xiyou Zhou xiyou@octoml.ai
Co-authored-by: Bohan Hou 32121147+spectrometerHBH@users.noreply.github.com
Co-authored-by: Siyuan Feng Hzfengsy@sjtu.edu.cn
Co-authored-by: Ruihang Lai lairuihangdongdong@qq.com
Co-authored-by: Wuwei Lin wuwei@apache.org
@junrushao1994 @Hzfengsy @comaniac
--
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] masahi commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791108913
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
I brought this up because I saw the upcoming code in https://github.com/junrushao1994/tvm/blob/meta-schedule/src/meta_schedule/schedule_rule/multi_level_tiling.cc where tensorcore stuff is hardcoded all over the place. I expect non-trivial refactoring before we can land this on `main`.
--
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 commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791113938
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
Well, you are right. We might want to remove all the tensor core hard-coding here and focusing on AutoScheduler alignment first during upstreaming
--
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 commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791080656
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
Yes. We want to support auto tensorization for all hardware platforms. use-tensor-core is probably not a good name here - maybe use-tensor-intrin could be a better one
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
Well, you are right. We might want to remove all the tensor core hard-coding here and focusing on AutoScheduler alignment first during upstreaming
--
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] masahi commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791072786
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
This is probably a broader question, but are we going to bake tensorcore specific logic into this core schedule rule? What if I want to support other matrix cores from Intel or AMD?
I feel like the design should be that backend specific rules be decoupled from the main "driver" that does actual tiling like the one introduced in this PR. cc @junrushao1994 @vinx13
--
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 commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
junrushao1994 commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791080656
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
Yes. We want to support auto tensorization for all hardware platforms. use-tensor-core is probably not a good name here - maybe use-tensor-intrin could be a better one
--
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] masahi commented on a change in pull request #10043: [MetaSchedule] Schedule Rule: Multi Level Tiling
Posted by GitBox <gi...@apache.org>.
masahi commented on a change in pull request #10043:
URL: https://github.com/apache/tvm/pull/10043#discussion_r791072786
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
This is probably a broader question, but are we going to bake tensorcore specific logic into this core schedule rule? What if I want to support other matrix cores from Intel or AMD?
I feel like the design should be that backend specific rules be decoupled from the main "driver" that does actual tiling like the one introduced in this PR. cc @junrushao1994 @vinx13
##########
File path: python/tvm/meta_schedule/schedule_rule/multi_level_tiling.py
##########
@@ -0,0 +1,88 @@
+# Licensed to the Apache Software Foundation (ASF) under one
+# or more contributor license agreements. See the NOTICE file
+# distributed with this work for additional information
+# regarding copyright ownership. The ASF licenses this file
+# to you under the Apache License, Version 2.0 (the
+# "License"); you may not use this file except in compliance
+# with the License. You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing,
+# software distributed under the License is distributed on an
+# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+# KIND, either express or implied. See the License for the
+# specific language governing permissions and limitations
+# under the License.
+"""Multi-level tiling with reuse."""
+from typing import Any, Dict, List, NamedTuple, Optional
+
+from tvm._ffi import register_object
+
+from .. import _ffi_api
+from .schedule_rule import ScheduleRule
+
+
+class ReuseType(NamedTuple):
+ """Reuse type."""
+
+ req: str
+ levels: List[int]
+ scope: str
+
+ def as_dict(self) -> Dict[str, Any]:
+ """Return the dict representation of the reuse type."""
+ return {
+ "req": self.req,
+ "levels": self.levels,
+ "scope": self.scope,
+ }
+
+
+@register_object("meta_schedule.MultiLevelTiling")
+class MultiLevelTiling(ScheduleRule):
+ """Multi-level tiling with reuse.
+
+ Parameters
+ ----------
+ structure : str
+ The tiling structure. Recommended:
+ - 'SSRSRS' on CPU
+ - 'SSSRRSRS' on GPU
+ tile_bind : Optional[List[str]]
+ For each level of tiles, which thread axis it is bound to. Recommended:
+ - None on CPU
+ - [blockIdx.x, vthread.x, threadIdx.x] on GPU
+ use_tensor_core : bool
+ Whether to apply tensor core wmma intrinsic for the computation
Review comment:
I brought this up because I saw the upcoming code in https://github.com/junrushao1994/tvm/blob/meta-schedule/src/meta_schedule/schedule_rule/multi_level_tiling.cc where tensorcore stuff is hardcoded all over the place. I expect non-trivial refactoring before we can land this on `main`.
--
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] Hzfengsy merged pull request #10043: [MetaSchedule][M4a] Schedule Rule: Multi-Level-Tiling
Posted by GitBox <gi...@apache.org>.
Hzfengsy merged pull request #10043:
URL: https://github.com/apache/tvm/pull/10043
--
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