You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "zxybazh (via GitHub)" <gi...@apache.org> on 2024/02/21 04:34:25 UTC

[I] [Bug] Tensorization Failure During Multilevel Tiling with Tensor Intrin [tvm]

zxybazh opened a new issue, #16614:
URL: https://github.com/apache/tvm/issues/16614

   ### Expected behavior
   
   MetaSchedule Tuning Works for the given Conv2d workload
   
   ### Actual behavior
   
   Triggers an error `ValueError: The block no longer exists in the IRModule` during application of schedule rule Multi-level tiling with tensor intrin. I notcied that `state->tensor_core_reindex_store` would point to a block that is already merged into another block via ComputeInline during application of `TileWithTensorIntrin`. 
   
   ### Environment
   
   Latest TVM Main
   
   ### Steps to reproduce
   
   ```
   import tvm
   from tvm.script import tir as T
   from tvm import meta_schedule as ms
   
   @T.prim_func(private=True)
   def fused_conv2d_add1(reshape3: T.Buffer((T.int64(50), T.int64(8), T.int64(72), T.int64(128)), "float16"), conv_in_weight: T.Buffer((T.int64(320), T.int64(8), T.int64(3), T.int64(3)), "float16"), lv23: T.Buffer((T.int64(1), T.int64(320), T.int64(1), T.int64(1)), "float16"), T_add_intermediate: T.Buffer((T.int64(50), T.int64(320), T.int64(72), T.int64(128)), "float16")):
       T.func_attr({"tir.noalias": T.bool(True)})
       # with T.block("root"):
       pad_temp = T.alloc_buffer((T.int64(50), T.int64(8), T.int64(74), T.int64(130)), "float16")
       conv2d_nchw_intermediate = T.alloc_buffer((T.int64(50), T.int64(320), T.int64(72), T.int64(128)), "float16")
       for i0, i1, i2, i3 in T.grid(T.int64(50), T.int64(8), T.int64(74), T.int64(130)):
           with T.block("pad_temp"):
               v_i0, v_i1, v_i2, v_i3 = T.axis.remap("SSSS", [i0, i1, i2, i3])
               T.reads(reshape3[v_i0, v_i1, v_i2 - T.int64(1), v_i3 - T.int64(1)])
               T.writes(pad_temp[v_i0, v_i1, v_i2, v_i3])
               pad_temp[v_i0, v_i1, v_i2, v_i3] = T.if_then_else(T.int64(1) <= v_i2 and v_i2 < T.int64(73) and T.int64(1) <= v_i3 and v_i3 < T.int64(129), reshape3[v_i0, v_i1, v_i2 - T.int64(1), v_i3 - T.int64(1)], T.float16(0))
       for nn, ff, yy, xx, rc, ry, rx in T.grid(T.int64(50), T.int64(320), T.int64(72), T.int64(128), T.int64(8), T.int64(3), T.int64(3)):
           with T.block("conv2d_nchw"):
               v_nn, v_ff, v_yy, v_xx, v_rc, v_ry, v_rx = T.axis.remap("SSSSRRR", [nn, ff, yy, xx, rc, ry, rx])
               T.reads(pad_temp[v_nn, v_rc, v_yy + v_ry, v_xx + v_rx], conv_in_weight[v_ff, v_rc, v_ry, v_rx])
               T.writes(conv2d_nchw_intermediate[v_nn, v_ff, v_yy, v_xx])
               with T.init():
                   conv2d_nchw_intermediate[v_nn, v_ff, v_yy, v_xx] = T.float16(0)
               conv2d_nchw_intermediate[v_nn, v_ff, v_yy, v_xx] = conv2d_nchw_intermediate[v_nn, v_ff, v_yy, v_xx] + pad_temp[v_nn, v_rc, v_yy + v_ry, v_xx + v_rx] * conv_in_weight[v_ff, v_rc, v_ry, v_rx]
       for ax0, ax1, ax2, ax3 in T.grid(T.int64(50), T.int64(320), T.int64(72), T.int64(128)):
           with T.block("T_add"):
               v_ax0, v_ax1, v_ax2, v_ax3 = T.axis.remap("SSSS", [ax0, ax1, ax2, ax3])
               T.reads(conv2d_nchw_intermediate[v_ax0, v_ax1, v_ax2, v_ax3], lv23[T.int64(0), v_ax1, T.int64(0), T.int64(0)])
               T.writes(T_add_intermediate[v_ax0, v_ax1, v_ax2, v_ax3])
               T_add_intermediate[v_ax0, v_ax1, v_ax2, v_ax3] = conv2d_nchw_intermediate[v_ax0, v_ax1, v_ax2, v_ax3] + lv23[T.int64(0), v_ax1, T.int64(0), T.int64(0)]
   
   target=tvm.target.Target("nvidia/nvidia-a10g")
   func = fused_conv2d_add1
   ms.tune_tir(func, target=target, max_trials_global=100, work_dir="./temp")
   ```
   


-- 
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.apache.org

For queries about this service, please contact Infrastructure at:
users@infra.apache.org


Re: [I] [Bug] Tensorization Failure During Multilevel Tiling with Tensor Intrin [tvm]

Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh closed issue #16614: [Bug] Tensorization Failure During Multilevel Tiling with Tensor Intrin
URL: https://github.com/apache/tvm/issues/16614


-- 
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: [I] [Bug] Tensorization Failure During Multilevel Tiling with Tensor Intrin [tvm]

Posted by "krishnab30 (via GitHub)" <gi...@apache.org>.
krishnab30 commented on issue #16614:
URL: https://github.com/apache/tvm/issues/16614#issuecomment-2020074885

   Hi @zxybazh , I am facing the same issue
   


-- 
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: [I] [Bug] Tensorization Failure During Multilevel Tiling with Tensor Intrin [tvm]

Posted by "zxybazh (via GitHub)" <gi...@apache.org>.
zxybazh commented on issue #16614:
URL: https://github.com/apache/tvm/issues/16614#issuecomment-2350218150

   Won't fix for now. Please feel free to reopen if this reproduces.


-- 
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