You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "FranklandJack (via GitHub)" <gi...@apache.org> on 2023/05/30 12:45:10 UTC
[GitHub] [tvm] FranklandJack opened a new pull request, #14981: [target] Use native architecture for `llvm` target
FranklandJack opened a new pull request, #14981:
URL: https://github.com/apache/tvm/pull/14981
Set the default `-device=` key for llvm targets based on the native architecture rather than hard coding to `cpu` which is x86 specific. This means that when llvm target triples are not specified we will test `arm_cpu` schedules on Arm®-based architectures and `cpu` schedules on x86 based architectures.
Fix any schedule test failures that result from this fix.
--
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
Re: [PR] [target] Use native architecture for `llvm` target [tvm]
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 commented on PR #14981:
URL: https://github.com/apache/tvm/pull/14981#issuecomment-2079018660
Closing as superseded by: https://github.com/apache/tvm/pull/16513
--
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] lhutton1 commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1230665774
##########
src/target/target_kind.cc:
##########
@@ -257,6 +257,12 @@ TargetJSON TestTargetParser(TargetJSON target) {
/********** Register Target kinds and attributes **********/
+#if defined(__arm__) || defined(__aarch64__)
+#define NATIVE_CPU "arm_cpu"
+#else
+#define NATIVE_CPU "cpu"
+#endif
Review Comment:
To make this more generic I wonder if we can use something like: https://github.com/apache/tvm/pull/14946/files#diff-f61b04b100f5145f2681340c81d3f2af221239594ed01e2e24896522329ce92cR2087. I agree that adding architectures iteratively will likely be simpler due to CI breakages, but this way we can set the precedence.
In addition, I think it would be good to add a warning when the host architecture could not be determined. E.g. something like: `Host architecture could not be determined. Compiling using 'cpu' fallback key which may result in unoptimized code. Please try adding a target key.` (probably needs thinking through a bit more) in order to prompt users.
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213409881
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -132,13 +132,16 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
plevel=15,
)
else:
+ # TODO(@FranklandJack)
+ # Investigate why this producing output tensor of
+ # incorrect dimensions in
+ # test_runtime_module_based_interface.py
Review Comment:
That's a really good point. Perhaps the best solution here is to just fix this schedule as part of this PR?
--
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
Re: [PR] [target] Use native architecture for `llvm` target [tvm]
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 closed pull request #14981: [target] Use native architecture for `llvm` target
URL: https://github.com/apache/tvm/pull/14981
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213102123
##########
src/target/target_kind.cc:
##########
@@ -257,6 +257,12 @@ TargetJSON TestTargetParser(TargetJSON target) {
/********** Register Target kinds and attributes **********/
+#if defined(__arm__) || defined(__aarch64__)
+#define NATIVE_CPU "arm_cpu"
+#else
+#define NATIVE_CPU "cpu"
+#endif
Review Comment:
A general solution would definitely be preferable, do you have a suggestion of how one would use the target parser to do this?
Since we've gone from being x86 specfic to arm and x86 specfic I'd argue this is still an improvement over what we had before. We can always add support for more targets in future PRs.
It's worth noting that this has exposed a lot of false positive tests on the `arm_cpu` target, if the same thing happens for other targets it might be advantageous to do this iteratively in order to keep the diff contained.
--
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] lhutton1 commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1230665774
##########
src/target/target_kind.cc:
##########
@@ -257,6 +257,12 @@ TargetJSON TestTargetParser(TargetJSON target) {
/********** Register Target kinds and attributes **********/
+#if defined(__arm__) || defined(__aarch64__)
+#define NATIVE_CPU "arm_cpu"
+#else
+#define NATIVE_CPU "cpu"
+#endif
Review Comment:
To make this more generic I wonder if we can use something like: https://github.com/apache/tvm/pull/14946/files#diff-f61b04b100f5145f2681340c81d3f2af221239594ed01e2e24896522329ce92cR2087. I agree that adding architectures iteratively will likely be simpler due to CI breakages, but this way we can set the precedence.
In addition, I think it would be good to add a warning when the host architecture could not be determined. E.g. something like: `Host architecture could not be determined. Compiling using 'cpu' fallback key which may result in unoptimized code. Please try adding a target.` (probably needs thinking through a bit more) in order to prompt users.
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1234934758
##########
python/tvm/topi/arm_cpu/injective.py:
##########
@@ -69,8 +69,10 @@ def schedule_injective(outs):
if list(s[x].op.axis):
# do not vectorize for broadcast
dtype = "uint16" if x.dtype == "bfloat16" else x.dtype
- (io, ii) = s[x].split(list(s[x].op.axis)[-1], 16 // np.dtype(dtype).itemsize)
- s[x].vectorize(ii)
+ # do not vectorize for custom data types
+ if 0 == dtype.count("custom"):
Review Comment:
Sure, fixed.
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213415889
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -284,8 +288,21 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
name="depthwise_conv2d_nchw.x86",
)
elif layout == "NHWC":
- assert kernel_layout == "HWOI"
- if target.features.has_asimd:
+ # TODO(@FranklandJack)
+ # Handle HWOI in arm_cpu scheudles/compute definition.
Review Comment:
Haha whoops! I should turn my spell checker on, will fix :)
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213417629
##########
tests/python/unittest/test_meta_schedule_relay_integration.py:
##########
@@ -420,6 +420,7 @@ def test_extract_task_arm_conv2d_nchwc():
def test_meta_schedule_te2primfunc_argument_order_and_lowering():
+ target = Target("llvm --num-cores=16")
Review Comment:
Ach, this is a debugging artefact that I forgot to clean up, will remove. Good catch!
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1226299048
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -132,13 +132,16 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
plevel=15,
)
else:
+ # TODO(@FranklandJack)
+ # Investigate why this producing output tensor of
+ # incorrect dimensions in
+ # test_runtime_module_based_interface.py
Review Comment:
For the time being I've just disabled the test that this causes to fail.
--
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] lhutton1 commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1230639408
##########
python/tvm/topi/arm_cpu/injective.py:
##########
@@ -69,8 +69,10 @@ def schedule_injective(outs):
if list(s[x].op.axis):
# do not vectorize for broadcast
dtype = "uint16" if x.dtype == "bfloat16" else x.dtype
- (io, ii) = s[x].split(list(s[x].op.axis)[-1], 16 // np.dtype(dtype).itemsize)
- s[x].vectorize(ii)
+ # do not vectorize for custom data types
+ if 0 == dtype.count("custom"):
Review Comment:
nit: this could just be `if dtype.count("custom")`
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -304,8 +321,11 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
and kernel.shape[3] == 1 # channel_multiplier == 1
and out_type.dtype == "int32"
and (
- (data.shape[3] % 4 == 0 and data.dtype == "int8" and target.features.has_dsp)
- or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ (
+ (data.shape[3] % 4 == 0 and data.dtype == "int8")
+ or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ )
+ and target.features.has_dsp
Review Comment:
The logic seems to have changed here, was this intentional?
--
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] lhutton1 commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "lhutton1 (via GitHub)" <gi...@apache.org>.
lhutton1 commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1234962491
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -304,8 +321,11 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
and kernel.shape[3] == 1 # channel_multiplier == 1
and out_type.dtype == "int32"
and (
- (data.shape[3] % 4 == 0 and data.dtype == "int8" and target.features.has_dsp)
- or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ (
+ (data.shape[3] % 4 == 0 and data.dtype == "int8")
+ or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ )
+ and target.features.has_dsp
Review Comment:
Aha yep agreed, seems like a bug :)
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1234939668
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -304,8 +321,11 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
and kernel.shape[3] == 1 # channel_multiplier == 1
and out_type.dtype == "int32"
and (
- (data.shape[3] % 4 == 0 and data.dtype == "int8" and target.features.has_dsp)
- or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ (
+ (data.shape[3] % 4 == 0 and data.dtype == "int8")
+ or (data.shape[3] % 2 == 0 and data.dtype == "int16")
+ )
+ and target.features.has_dsp
Review Comment:
I think so, I think the idea was that we can use the dsp schedules on int8 and int16 only where it is supported, previously it would have used the dsp schedule for int16 regardless.
--
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] ekalda commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "ekalda (via GitHub)" <gi...@apache.org>.
ekalda commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1211974551
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -132,13 +132,16 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
plevel=15,
)
else:
+ # TODO(@FranklandJack)
+ # Investigate why this producing output tensor of
+ # incorrect dimensions in
+ # test_runtime_module_based_interface.py
Review Comment:
Are there any other tests that exercise the same schedule that pass? I'm quite worried about commenting out an Arm specific schedule, essentially defaulting to use x86 schedule for this data and kernel layout.
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -284,8 +288,21 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
name="depthwise_conv2d_nchw.x86",
)
elif layout == "NHWC":
- assert kernel_layout == "HWOI"
- if target.features.has_asimd:
+ # TODO(@FranklandJack)
+ # Handle HWOI in arm_cpu scheudles/compute definition.
+ if kernel_layout == "HWOI":
+ logger.warning(
+ """depthwise_conv2d with layout NHWC and HWOI
+ kernel layout is not optimized for arm_cpu target.
+ """
+ )
+ strategy.add_implementation(
+ wrap_compute_conv2d(topi.nn.depthwise_conv2d_nhwc, need_kernel_layout=True),
+ wrap_topi_schedule(conv2d_generic.schedule_depthwise_conv2d_nhwc),
+ name="depthwise_conv2d_nhwc.generic",
+ )
+
+ elif target.features.has_asimd:
Review Comment:
It looks like the logic is changing quite a bit there... So if we have NHWC with HWOI, we add first the generic schedule, then if there is vector extension, we also add the arm specific schedule? Why is that? Also the DSP schedule below requires HWOI, but doesn't check for it explicitly since it is relying on the assert that is being removed.
##########
python/tvm/topi/arm_cpu/conv2d.py:
##########
@@ -509,3 +510,24 @@ def conv2d_nhwc_dsp(cfg, data, kernel, strides, padding, dilation, out_dtype):
def schedule_conv2d_nhwc_dsp(cfg, outs):
"""Create schedule for conv2d_nhwc_dsp"""
return conv2d_nhwc_dsp_schedule(cfg, outs)
+
+
+@conv2d_infer_layout.register("arm_cpu")
+def _conv2d_infer_layout(workload, cfg):
+ _, data, kernel, strides, padding, dilation, layout, _, dtype = workload
+ batch_size, in_channel, in_height, in_width = data[1]
+ out_channel, _, k_height, k_width = kernel[1]
+ idxdiv = tvm.tir.indexdiv
+
+ pt, pl, pb, pr = get_pad_tuple(padding, (k_height, k_width))
+ hdilation, wdilation = dilation if isinstance(dilation, (tuple, list)) else (dilation, dilation)
+ dilated_kernel_h = (k_height - 1) * hdilation + 1
+ dilated_kernel_w = (k_width - 1) * wdilation + 1
+ out_height = idxdiv(in_height + pt + pb - dilated_kernel_h, strides[0]) + 1
+ out_width = idxdiv(in_width + pl + pr - dilated_kernel_w, strides[1]) + 1
+ tile_ic, tile_oc = cfg["tile_ic"].size[-1], cfg["tile_oc"].size[-1]
+ in_shape = (batch_size, idxdiv(in_channel, tile_ic), in_height, in_width, tile_ic)
+ in_layout = f"NCHW{tile_ic}c"
+ out_shape = (batch_size, idxdiv(out_channel, tile_oc), out_height, out_width, tile_oc)
+ out_layout = f"NCHW{tile_oc}c"
+ return ((in_shape, in_layout),), ((out_shape, out_layout),)
Review Comment:
Is this function identical to the ones in `x86` and `intel_graphics`? If yes, then while you are at it, maybe move the body of that function to some common location (e.g. `python/tvm/topi/utils.py`) to reduce the amount of duplicated code.
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -284,8 +288,21 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
name="depthwise_conv2d_nchw.x86",
)
elif layout == "NHWC":
- assert kernel_layout == "HWOI"
- if target.features.has_asimd:
+ # TODO(@FranklandJack)
+ # Handle HWOI in arm_cpu scheudles/compute definition.
Review Comment:
Nit:
```suggestion
# Handle HWOI in arm_cpu schedules/compute definition.
```
##########
tests/python/unittest/test_meta_schedule_relay_integration.py:
##########
@@ -420,6 +420,7 @@ def test_extract_task_arm_conv2d_nchwc():
def test_meta_schedule_te2primfunc_argument_order_and_lowering():
+ target = Target("llvm --num-cores=16")
Review Comment:
Just curious, what does that change do?
--
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] kparzysz-quic commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "kparzysz-quic (via GitHub)" <gi...@apache.org>.
kparzysz-quic commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213068645
##########
src/target/target_kind.cc:
##########
@@ -257,6 +257,12 @@ TargetJSON TestTargetParser(TargetJSON target) {
/********** Register Target kinds and attributes **********/
+#if defined(__arm__) || defined(__aarch64__)
+#define NATIVE_CPU "arm_cpu"
+#else
+#define NATIVE_CPU "cpu"
+#endif
Review Comment:
This isn't really much of an improvement over being "x86-specific": now the code is "x86 and ARM-specific". Isn't it possible to set keys using target parser? It would be nice to have a more general solution for 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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1213425076
##########
python/tvm/topi/arm_cpu/conv2d.py:
##########
@@ -509,3 +510,24 @@ def conv2d_nhwc_dsp(cfg, data, kernel, strides, padding, dilation, out_dtype):
def schedule_conv2d_nhwc_dsp(cfg, outs):
"""Create schedule for conv2d_nhwc_dsp"""
return conv2d_nhwc_dsp_schedule(cfg, outs)
+
+
+@conv2d_infer_layout.register("arm_cpu")
+def _conv2d_infer_layout(workload, cfg):
+ _, data, kernel, strides, padding, dilation, layout, _, dtype = workload
+ batch_size, in_channel, in_height, in_width = data[1]
+ out_channel, _, k_height, k_width = kernel[1]
+ idxdiv = tvm.tir.indexdiv
+
+ pt, pl, pb, pr = get_pad_tuple(padding, (k_height, k_width))
+ hdilation, wdilation = dilation if isinstance(dilation, (tuple, list)) else (dilation, dilation)
+ dilated_kernel_h = (k_height - 1) * hdilation + 1
+ dilated_kernel_w = (k_width - 1) * wdilation + 1
+ out_height = idxdiv(in_height + pt + pb - dilated_kernel_h, strides[0]) + 1
+ out_width = idxdiv(in_width + pl + pr - dilated_kernel_w, strides[1]) + 1
+ tile_ic, tile_oc = cfg["tile_ic"].size[-1], cfg["tile_oc"].size[-1]
+ in_shape = (batch_size, idxdiv(in_channel, tile_ic), in_height, in_width, tile_ic)
+ in_layout = f"NCHW{tile_ic}c"
+ out_shape = (batch_size, idxdiv(out_channel, tile_oc), out_height, out_width, tile_oc)
+ out_layout = f"NCHW{tile_oc}c"
+ return ((in_shape, in_layout),), ((out_shape, out_layout),)
Review Comment:
Yep it was copied directly from there. I guess this will only work for scheduling definitions that include 'tile_ic' and 'tile_oc' parameters (which is true for x86, intel_graphics and arm_cpu) but maybe not others? Is that okay? I guess we can document that.
--
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] FranklandJack commented on a diff in pull request #14981: [target] Use native architecture for `llvm` target
Posted by "FranklandJack (via GitHub)" <gi...@apache.org>.
FranklandJack commented on code in PR #14981:
URL: https://github.com/apache/tvm/pull/14981#discussion_r1214174551
##########
python/tvm/relay/op/strategy/arm_cpu.py:
##########
@@ -284,8 +288,21 @@ def conv2d_strategy_arm_cpu(attrs, inputs, out_type, target):
name="depthwise_conv2d_nchw.x86",
)
elif layout == "NHWC":
- assert kernel_layout == "HWOI"
- if target.features.has_asimd:
+ # TODO(@FranklandJack)
+ # Handle HWOI in arm_cpu scheudles/compute definition.
+ if kernel_layout == "HWOI":
+ logger.warning(
+ """depthwise_conv2d with layout NHWC and HWOI
+ kernel layout is not optimized for arm_cpu target.
+ """
+ )
+ strategy.add_implementation(
+ wrap_compute_conv2d(topi.nn.depthwise_conv2d_nhwc, need_kernel_layout=True),
+ wrap_topi_schedule(conv2d_generic.schedule_depthwise_conv2d_nhwc),
+ name="depthwise_conv2d_nhwc.generic",
+ )
+
+ elif target.features.has_asimd:
Review Comment:
Ah, so I think this might be a bug. The condition I think should be `if kernel_layout != "HWOI"` so that we fallback on the generic condition in the unsupported case, does that make sense?
--
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] tvm-bot commented on pull request #14981: [target] Use native architecture for `llvm` target
Posted by "tvm-bot (via GitHub)" <gi...@apache.org>.
tvm-bot commented on PR #14981:
URL: https://github.com/apache/tvm/pull/14981#issuecomment-1568371635
<!---bot-comment-->
Thanks for contributing to TVM! Please refer to the contributing guidelines https://tvm.apache.org/docs/contribute/ for useful information and tips. Please request code reviews from [Reviewers](https://github.com/apache/incubator-tvm/blob/master/CONTRIBUTORS.md#reviewers) by @-ing them in a comment.
<!--bot-comment-ccs-start-->
* No users to tag found in teams: `target` <sub>See [#10317](https://github.com/apache/tvm/issues/10317) for details</sub><!--bot-comment-ccs-end-->
<sub>Generated by [tvm-bot](https://github.com/apache/tvm/blob/main/ci/README.md#github-actions)</sub>
--
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