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 2021/09/30 10:27:02 UTC
[GitHub] [tvm] mbaret opened a new pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
mbaret opened a new pull request #9159:
URL: https://github.com/apache/tvm/pull/9159
This patch addresses the review comments left on https://github.com/apache/tvm/pull/8806. Small changes are made to the API of ScheduleBuilder so that it can be reused in the NPU compiler, allowing us to remove a largely duplicated piece of code.
--
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] mbs-octoml commented on a change in pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
mbs-octoml commented on a change in pull request #9159:
URL: https://github.com/apache/tvm/pull/9159#discussion_r722359163
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir/scheduler.py
##########
@@ -15,17 +15,17 @@
# specific language governing permissions and limitations
# under the License.
# pylint: disable=invalid-name, unused-argument
-"""Different schedulers for Arm(R) Ethos(TM)-U NPU"""
+"""Scheduling for Arm(R) Ethos(TM)-U NPU."""
import tvm
-def schedule(te_graph, const_dict, cascader=None):
- """Schedule a TE graph for NPU compilation.
+def schedule(cached_func, const_dict, cascader=None):
+ """Schedule a CachedFunc for NPU compilation.
Parameters
----------
- te_graph
- The TE graph to schedule.
+ cached_func : CachedFunc
Review comment:
Heads up we're trying to move both the TECompiler and CachedFunc structures to be internal only. I think the idea is all the cross-reference stuff accumulated in CachedFunc would be captured as 'official' attributes on the call or defn. So obviously we'll need to include this in that refactor.
--
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] mbaret commented on a change in pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
mbaret commented on a change in pull request #9159:
URL: https://github.com/apache/tvm/pull/9159#discussion_r725158741
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -447,81 +455,75 @@ def _create_npu_feature_map(serial_feature_map):
)
nfm.layout = layout_map[layout]
nfm.strides = vapi.NpuShape3D(
- int(serial_feature_map.stride_h.value),
- int(serial_feature_map.stride_w.value),
- int(serial_feature_map.stride_c.value),
+ int(serial_feature_map.stride_h),
+ int(serial_feature_map.stride_w),
+ int(serial_feature_map.stride_c),
)
return nfm
-def _create_npu_kernel(serial_kernel):
+def _create_npu_kernel(serial_kernel: spec.SerialKernel) -> vapi.NpuKernel:
"""This is a helper function to capture a list
- of arguments to create Vela NpuKernel object
+ of arguments to create Vela NpuKernel object.
"""
nknl = vapi.NpuKernel(
- w=int(serial_kernel.width.value),
- h=int(serial_kernel.height.value),
- stride_x=int(serial_kernel.stride_w.value),
- stride_y=int(serial_kernel.stride_h.value),
- dilation_x=int(serial_kernel.dilation_w.value),
- dilation_y=int(serial_kernel.dilation_h.value),
+ w=int(serial_kernel.width),
+ h=int(serial_kernel.height),
+ stride_x=int(serial_kernel.stride_w),
+ stride_y=int(serial_kernel.stride_h),
+ dilation_x=int(serial_kernel.dilation_w),
+ dilation_y=int(serial_kernel.dilation_h),
)
return nknl
-def _create_npu_address_range(serial_address_range):
+def _create_npu_address_range(
+ serial_address_range: spec.SerialAddressRange,
+) -> vapi.NpuAddressRange:
"""This is a helper function to capture a list
- of arguments to create Vela NpuAddressRange object
+ of arguments to create Vela NpuAddressRange object.
"""
addr_range = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_address_range.address,
- length=int(serial_address_range.length.value),
+ length=int(serial_address_range.length),
)
return addr_range
def _create_npu_quantization(
- scale,
- zero_point,
-):
+ scale: Union[tvm.tir.FloatImm, float],
+ zero_point: Union[tvm.tir.IntImm, int],
+) -> vapi.NpuQuantization:
"""This is a helper function to capture a list
- of arguments to create Vela NpuQuantization object
+ of arguments to create Vela NpuQuantization object.
"""
- # Scale could be an ndarray if per-channel quantization is available
- if not isinstance(scale, tvm.tir.expr.Load):
- if isinstance(scale.value, float):
- scale = np.single(scale.value)
- else:
- assert isinstance(scale.value.value, float)
- scale = np.single(scale.value.value)
- q_params = vapi.NpuQuantization(scale_f32=scale, zero_point=zero_point.value)
- return q_params
+ return vapi.NpuQuantization(scale_f32=float(scale), zero_point=int(zero_point))
Review comment:
This function was doing something it didn't need to (per-channel scales never actually appear here), so I just simplified it to remove that confusion. I'll amend the commit message to include reference to this.
--
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] mbaret commented on pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
mbaret commented on pull request #9159:
URL: https://github.com/apache/tvm/pull/9159#issuecomment-938815404
Apologies for the delay fixing CI here, having some environment problems locally. Will hopefully be fixed by next week.
--
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] areusch merged pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
areusch merged pull request #9159:
URL: https://github.com/apache/tvm/pull/9159
--
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] areusch commented on pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
areusch commented on pull request #9159:
URL: https://github.com/apache/tvm/pull/9159#issuecomment-933751071
cc @mbs-octoml @jroesch PTAL
--
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] areusch merged pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
areusch merged pull request #9159:
URL: https://github.com/apache/tvm/pull/9159
--
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] mbs-octoml commented on a change in pull request #9159: Address review comments on Arm(R) Ethos(TM)-U PR 3/6
Posted by GitBox <gi...@apache.org>.
mbs-octoml commented on a change in pull request #9159:
URL: https://github.com/apache/tvm/pull/9159#discussion_r721674546
##########
File path: python/tvm/relay/backend/contrib/ethosu/tir_to_cs_translator.py
##########
@@ -447,81 +455,75 @@ def _create_npu_feature_map(serial_feature_map):
)
nfm.layout = layout_map[layout]
nfm.strides = vapi.NpuShape3D(
- int(serial_feature_map.stride_h.value),
- int(serial_feature_map.stride_w.value),
- int(serial_feature_map.stride_c.value),
+ int(serial_feature_map.stride_h),
+ int(serial_feature_map.stride_w),
+ int(serial_feature_map.stride_c),
)
return nfm
-def _create_npu_kernel(serial_kernel):
+def _create_npu_kernel(serial_kernel: spec.SerialKernel) -> vapi.NpuKernel:
"""This is a helper function to capture a list
- of arguments to create Vela NpuKernel object
+ of arguments to create Vela NpuKernel object.
"""
nknl = vapi.NpuKernel(
- w=int(serial_kernel.width.value),
- h=int(serial_kernel.height.value),
- stride_x=int(serial_kernel.stride_w.value),
- stride_y=int(serial_kernel.stride_h.value),
- dilation_x=int(serial_kernel.dilation_w.value),
- dilation_y=int(serial_kernel.dilation_h.value),
+ w=int(serial_kernel.width),
+ h=int(serial_kernel.height),
+ stride_x=int(serial_kernel.stride_w),
+ stride_y=int(serial_kernel.stride_h),
+ dilation_x=int(serial_kernel.dilation_w),
+ dilation_y=int(serial_kernel.dilation_h),
)
return nknl
-def _create_npu_address_range(serial_address_range):
+def _create_npu_address_range(
+ serial_address_range: spec.SerialAddressRange,
+) -> vapi.NpuAddressRange:
"""This is a helper function to capture a list
- of arguments to create Vela NpuAddressRange object
+ of arguments to create Vela NpuAddressRange object.
"""
addr_range = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_address_range.address,
- length=int(serial_address_range.length.value),
+ length=int(serial_address_range.length),
)
return addr_range
def _create_npu_quantization(
- scale,
- zero_point,
-):
+ scale: Union[tvm.tir.FloatImm, float],
+ zero_point: Union[tvm.tir.IntImm, int],
+) -> vapi.NpuQuantization:
"""This is a helper function to capture a list
- of arguments to create Vela NpuQuantization object
+ of arguments to create Vela NpuQuantization object.
"""
- # Scale could be an ndarray if per-channel quantization is available
- if not isinstance(scale, tvm.tir.expr.Load):
- if isinstance(scale.value, float):
- scale = np.single(scale.value)
- else:
- assert isinstance(scale.value.value, float)
- scale = np.single(scale.value.value)
- q_params = vapi.NpuQuantization(scale_f32=scale, zero_point=zero_point.value)
- return q_params
+ return vapi.NpuQuantization(scale_f32=float(scale), zero_point=int(zero_point))
Review comment:
The only semantic change, right? Can you explain it in the PR summary, thx.
--
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