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 2020/05/27 16:22:37 UTC
[GitHub] [incubator-tvm] cchung100m opened a new pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
cchung100m opened a new pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684
Hi @kevinthesun
Following the issue #5215, I would appreciate that if you can help to review this PR, thank you.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] cchung100m commented on pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
cchung100m commented on pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#issuecomment-636443296
Hi @kevinthesun
Thanks for the prompt reply and reviews. :)
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431502771
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -142,13 +142,14 @@ def _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec):
s[data_vec].unroll(vw)
if isinstance(kernel_vec.op, tvm.te.ComputeOp) and kernel_vec.name == 'kernel_vec':
+ co, ci, kh, kw, vc = s[kernel_vec].op.axis
Review comment:
Looks like this part share the logic with arm_cpu and mali. We can move this into [conv2d_nchw_spatial_pack compute](https://github.com/apache/incubator-tvm/blob/master/topi/python/topi/arm_cpu/conv2d_spatial_pack.py#L27). An example is in x86 conv2d: https://github.com/apache/incubator-tvm/blob/master/topi/python/topi/x86/conv2d.py#L188-L201
With this method, we can fix mali/arm_cpu/bifrost in one shot.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431505903
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -142,13 +142,14 @@ def _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec):
s[data_vec].unroll(vw)
if isinstance(kernel_vec.op, tvm.te.ComputeOp) and kernel_vec.name == 'kernel_vec':
+ co, ci, kh, kw, vc = s[kernel_vec].op.axis
Review comment:
Actually I'm not quite sure whether we can directly do this kind of placeholder replacement inside schedule. https://github.com/apache/incubator-tvm/pull/5511 for mali might also has issue. Put placeholder replacement in compute would be more convenient.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431502771
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -142,13 +142,14 @@ def _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec):
s[data_vec].unroll(vw)
if isinstance(kernel_vec.op, tvm.te.ComputeOp) and kernel_vec.name == 'kernel_vec':
+ co, ci, kh, kw, vc = s[kernel_vec].op.axis
Review comment:
Looks like this part shares the logic with arm_cpu and mali. We can move this into [conv2d_nchw_spatial_pack compute](https://github.com/apache/incubator-tvm/blob/master/topi/python/topi/arm_cpu/conv2d_spatial_pack.py#L27). An example is in x86 conv2d: https://github.com/apache/incubator-tvm/blob/master/topi/python/topi/x86/conv2d.py#L188-L201
With this method, we can fix mali/arm_cpu/bifrost in one shot.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431504200
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -370,12 +371,7 @@ def _schedule_winograd(cfg, s, op):
s[G].compute_inline()
eps, _, _, _ = s[U].op.axis
y, _, _, _ = s[padded_kernel].op.axis
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # Kernel transformation will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[U].pragma(eps, 'debug_skip_region')
- s[padded_kernel].pragma(y, 'debug_skip_region')
- else:
+ if not autotvm.GLOBAL_SCOPE.in_tuning:
Review comment:
Also winograd kernel transformation is different from spatial_pack. We need to have a different path inside winograd compute to generate placeholder.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431505903
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -142,13 +142,14 @@ def _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec):
s[data_vec].unroll(vw)
if isinstance(kernel_vec.op, tvm.te.ComputeOp) and kernel_vec.name == 'kernel_vec':
+ co, ci, kh, kw, vc = s[kernel_vec].op.axis
Review comment:
Actually I'm not quite sure whether we can directly do this kind of placeholder replacement inside schedule. https://github.com/apache/incubator-tvm/pull/5511 for mali might also has issue. Putting placeholder replacement in compute would be more convenient.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#issuecomment-636517489
Thanks @cchung100m
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] cchung100m commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
cchung100m commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r432926021
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -370,12 +371,7 @@ def _schedule_winograd(cfg, s, op):
s[G].compute_inline()
eps, _, _, _ = s[U].op.axis
y, _, _, _ = s[padded_kernel].op.axis
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # Kernel transformation will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[U].pragma(eps, 'debug_skip_region')
- s[padded_kernel].pragma(y, 'debug_skip_region')
- else:
+ if not autotvm.GLOBAL_SCOPE.in_tuning:
Review comment:
Yes. I put the placeholder in `_decl_winograd` to biforst/con2d.py and mali/con2d.py
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r431504200
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -370,12 +371,7 @@ def _schedule_winograd(cfg, s, op):
s[G].compute_inline()
eps, _, _, _ = s[U].op.axis
y, _, _, _ = s[padded_kernel].op.axis
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # Kernel transformation will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[U].pragma(eps, 'debug_skip_region')
- s[padded_kernel].pragma(y, 'debug_skip_region')
- else:
+ if not autotvm.GLOBAL_SCOPE.in_tuning:
Review comment:
Also winograd kernel transformation is different from spatial_pack. We need to have a different path to generate placeholder.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#issuecomment-636407068
Also can you fix winograd kernel replacement?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] cchung100m commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
cchung100m commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r432817217
##########
File path: topi/python/topi/bifrost/conv2d.py
##########
@@ -142,13 +142,14 @@ def _schedule_spatial_pack(cfg, s, output, conv, data_vec, kernel_vec):
s[data_vec].unroll(vw)
if isinstance(kernel_vec.op, tvm.te.ComputeOp) and kernel_vec.name == 'kernel_vec':
+ co, ci, kh, kw, vc = s[kernel_vec].op.axis
Review comment:
Hi @kevinthesun
Thanks for the review, I move the placeholder replacement to compute.
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun merged pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun merged pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
[GitHub] [incubator-tvm] kevinthesun commented on a change in pull request #5684: [AutoTVM][TOPI] Fix bifrost spatial packing conv2d auto tune
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5684:
URL: https://github.com/apache/incubator-tvm/pull/5684#discussion_r432900758
##########
File path: topi/python/topi/arm_cpu/conv2d_spatial_pack.py
##########
@@ -267,9 +270,13 @@ def conv2d_spatial_pack_nhwc(cfg, data, kernel, strides, padding, dilation, out_
data_vec = te.compute(dvshape, lambda n, oho, owo, ohi, owi, ic:
data_pad[n][oho*OHI*HSTR+ohi][owo*OWI*WSTR+owi][ic],
name='data_vec')
- kernel_vec = te.compute(kvshape, lambda oco, kh, kw, ic, oci: \
- kernel[kh][kw][ic][oco*OCI+oci],
- name='kernel_vec')
+
+ if autotvm.GLOBAL_SCOPE.in_tuning:
Review comment:
Change schedule for arm_cpu conv2d as well?
----------------------------------------------------------------
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.
For queries about this service, please contact Infrastructure at:
users@infra.apache.org