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/05 15:58:39 UTC
[GitHub] [incubator-tvm] cchung100m opened a new pull request #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
cchung100m opened a new pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511
Hi @FrozenGene @kevinthesun @comaniac
I am interested in issue #5215, however, I am not familiar with `TOPI` and `backend`. I would appreciate that if you can guide me to complete this work, thank you in advance. :)
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#issuecomment-634164849
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] kevinthesun commented on a change in pull request #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r421056807
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -138,20 +138,15 @@ 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':
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # kernel packing will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[kernel_vec].pragma(s[kernel_vec].op.axis[0], 'debug_skip_region')
- else:
- max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
- co, ci, kh, kw, vc = s[kernel_vec].op.axis
- fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
- fused, vec = s[kernel_vec].split(fused, VC)
- bb, tt = s[kernel_vec].split(fused, max_threads)
- s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
- s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
- if VC in vec_size:
- s[kernel_vec].vectorize(vec)
+ max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
Review comment:
Thank you for working on this. You also need to replace kernel with the correct layout placeholder in compute when autotvm.GLOBAL_SCOPE.in_tuning is True. Look at https://github.com/apache/incubator-tvm/pull/5200/files#diff-7d0b52f92ecc0d246c17fa17f90d2a8fR189
----------------------------------------------------------------
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] tqchen commented on pull request #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
tqchen commented on pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#issuecomment-629505570
@kevinthesun please followup
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
cchung100m commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r422642720
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -138,20 +138,15 @@ 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':
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # kernel packing will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[kernel_vec].pragma(s[kernel_vec].op.axis[0], 'debug_skip_region')
- else:
- max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
- co, ci, kh, kw, vc = s[kernel_vec].op.axis
- fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
- fused, vec = s[kernel_vec].split(fused, VC)
- bb, tt = s[kernel_vec].split(fused, max_threads)
- s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
- s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
- if VC in vec_size:
- s[kernel_vec].vectorize(vec)
+ max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
Review comment:
Hi @kevinthesun
Thanks for the kind explanation.
Should we keep removing the `autotvm.GLOBAL_SCOPE.in_tuning:` block here, and replace kernel with the correct layout placeholder in below position?
https://github.com/apache/incubator-tvm/blob/418da931c1806b35266bc7e12c3c8ffb01fa034b/topi/python/topi/mali/conv2d.py#L456
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r426125601
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -138,20 +138,15 @@ 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':
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # kernel packing will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[kernel_vec].pragma(s[kernel_vec].op.axis[0], 'debug_skip_region')
- else:
- max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
- co, ci, kh, kw, vc = s[kernel_vec].op.axis
- fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
- fused, vec = s[kernel_vec].split(fused, VC)
- bb, tt = s[kernel_vec].split(fused, max_threads)
- s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
- s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
- if VC in vec_size:
- s[kernel_vec].vectorize(vec)
+ max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
Review comment:
We need to do the replacement in 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] cchung100m commented on pull request #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
cchung100m commented on pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#issuecomment-629620151
Hi @kevinthesun
Thanks for the feedback. I am still thinking how to do this task, currently. Therefore, I will turn it to draft now.
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun merged pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
cchung100m commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r430330419
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -345,16 +347,11 @@ def _schedule_winograd(cfg, s, op):
kernel, G = s[U].op.input_tensors
s[G].compute_inline()
eps, nu, co, ci, vco, = s[U].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')
- else:
- r_kh, r_kw = s[U].op.reduce_axis
- s[U].reorder(co, ci, eps, nu, r_kh, r_kw, vco)
- _ = [s[U].unroll(x) for x in [eps, nu, r_kh, r_kw]]
- s[U].vectorize(vco)
- tile_and_bind(s, U, co, ci, 1, 256)
+ r_kh, r_kw = s[U].op.reduce_axis
Review comment:
Hi @kevinthesun
Thanks for the review. I update the part you mentioned.
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r430091383
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -345,16 +347,11 @@ def _schedule_winograd(cfg, s, op):
kernel, G = s[U].op.input_tensors
s[G].compute_inline()
eps, nu, co, ci, vco, = s[U].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')
- else:
- r_kh, r_kw = s[U].op.reduce_axis
- s[U].reorder(co, ci, eps, nu, r_kh, r_kw, vco)
- _ = [s[U].unroll(x) for x in [eps, nu, r_kh, r_kw]]
- s[U].vectorize(vco)
- tile_and_bind(s, U, co, ci, 1, 256)
+ r_kh, r_kw = s[U].op.reduce_axis
Review comment:
This part should not be required in autotuning?
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
cchung100m commented on pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#issuecomment-633621823
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
cchung100m commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r421266463
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -138,20 +138,15 @@ 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':
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # kernel packing will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[kernel_vec].pragma(s[kernel_vec].op.axis[0], 'debug_skip_region')
- else:
- max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
- co, ci, kh, kw, vc = s[kernel_vec].op.axis
- fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
- fused, vec = s[kernel_vec].split(fused, VC)
- bb, tt = s[kernel_vec].split(fused, max_threads)
- s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
- s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
- if VC in vec_size:
- s[kernel_vec].vectorize(vec)
+ max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
Review comment:
Hi @kevinthesun
Thanks for your kind feedback. I thought it is a schedule of spatial packing for conv2d, so that, I remove the `autotvm.GLOBAL_SCOPE.in_tuning` here. Just like: https://github.com/apache/incubator-tvm/pull/5200/files#diff-1646331f629f3c6af025167c79d09427L98
Can you elaborate a bit about what do you mean?
many thanks,
----------------------------------------------------------------
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 #5511: [AutoTVM][TOPI] AutoTVM incorrect measurement
Posted by GitBox <gi...@apache.org>.
kevinthesun commented on a change in pull request #5511:
URL: https://github.com/apache/incubator-tvm/pull/5511#discussion_r422409992
##########
File path: topi/python/topi/mali/conv2d.py
##########
@@ -138,20 +138,15 @@ 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':
- if autotvm.GLOBAL_SCOPE.in_tuning:
- # kernel packing will be pre-computed during compilation, so we skip
- # this part to make tuning records correct
- s[kernel_vec].pragma(s[kernel_vec].op.axis[0], 'debug_skip_region')
- else:
- max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
- co, ci, kh, kw, vc = s[kernel_vec].op.axis
- fused = s[kernel_vec].fuse(co, ci, kh, kw, vc)
- fused, vec = s[kernel_vec].split(fused, VC)
- bb, tt = s[kernel_vec].split(fused, max_threads)
- s[kernel_vec].bind(bb, te.thread_axis("blockIdx.x"))
- s[kernel_vec].bind(tt, te.thread_axis("threadIdx.x"))
- if VC in vec_size:
- s[kernel_vec].vectorize(vec)
+ max_threads = tvm.target.Target.current(allow_none=False).max_num_threads
Review comment:
While doing autotvm, all the workloads fetched are in the original data/kernel layouts, for example NCHW/OIHW. It means we need to do layout conversion first. For spatial_pack this is done inside compute. However, we need to skip this layout conversion stage while doing autotuning. Previously we use debug_skip_region which will cause inaccurate measurement. Another way to skip it is to change the input kernel tensor to a new placeholder with converted layout.
----------------------------------------------------------------
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