[Bug] InternalError: Check failed: (ret_ex.dtype() == var.dtype()) is false

### Actual behavior

```
Traceback (most recent call last):
File "/data/qshenaf/remote_pc/TirFuzz/bugs/05-03_20-50/topi.nn.group_norm_4.py", line 20, in
tvm.build(opt_mod, target='llvm')
File "/data/qshenaf/envs/tvm/python/tvm/driver/build_module.py", line 59, in build
return tvm.tir.build(mod, target, pipeline)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/tir/build.py", line 173, in build
mod = pipeline(mod)
^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/ir/transform.py", line 238, in __call__
return _ffi_transform_api.RunPass(self, mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "tvm/_ffi/_cython/./packed_func.pxi", line 339, in tvm._ffi._cy3.core.PackedFuncBase.__call__
File "tvm/_ffi/_cython/./packed_func.pxi", line 270, in tvm._ffi._cy3.core.FuncCall
File "tvm/_ffi/_cython/./packed_func.pxi", line 259, in tvm._ffi._cy3.core.FuncCall3
File "tvm/_ffi/_cython/./base.pxi", line 185, in tvm._ffi._cy3.core.CHECK_CALL
File "/data/qshenaf/envs/tvm/python/tvm/_ffi/base.py", line 468, in raise_last_ffi_error
raise py_err
File "tvm/_ffi/_cython/./packed_func.pxi", line 56, in tvm._ffi._cy3.core.tvm_callback
File "/data/qshenaf/envs/tvm/python/tvm/tir/pipeline.py", line 122, in _pipeline
mod = tvm.ir.transform.Sequential(passes)(mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/python/tvm/ir/transform.py", line 238, in __call__
return _ffi_transform_api.RunPass(self, mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "tvm/_ffi/_cython/./packed_func.pxi", line 339, in tvm._ffi._cy3.core.PackedFuncBase.__call__
File "tvm/_ffi/_cython/./packed_func.pxi", line 270, in tvm._ffi._cy3.core.FuncCall
File "tvm/_ffi/_cython/./packed_func.pxi", line 259, in tvm._ffi._cy3.core.FuncCall3
File "tvm/_ffi/_cython/./base.pxi", line 185, in tvm._ffi._cy3.core.CHECK_CALL
File "/data/qshenaf/envs/tvm/src/tir/ir/transform.cc", line 121, in tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
func = pass_func(std::move(func), mod, pass_ctx);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 255, in operator()
return PlanAndUpdateBufferAllocationLocation(std::move(f));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 247, in tvm::tir::PlanAndUpdateBufferAllocationLocation(tvm::tir::PrimFunc)
fptr->body = locator(fptr->body);
^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 178, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::BlockNode const*)
Stmt stmt = StmtMutator::VisitStmt_(op);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 211, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)
return MutateArray(self, arr, fmutate);
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 184, in tvm::runtime::Array::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Arrayconst&)::{lambda(tvm::tir::Stmt const&)#1})
Array copy = arr.Map(fmutate);
^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 652, in tvm::runtime::Array::value, void>::type> tvm::runtime::Array::Map const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}) const
return Array(MapHelper(data_, fmap));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 877, in tvm::runtime::ObjectPtr tvm::runtime::Array::MapHelper const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::runtime::ObjectPtr, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
U mapped = fmap(DowncastNoCheck(*it));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 210, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
auto fmutate = [self](const Stmt& s) { return self->VisitStmt(s); };
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 145, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
auto node = Downcast(StmtMutator::VisitStmt_(op));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 211, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)
return MutateArray(self, arr, fmutate);
^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 184, in tvm::runtime::Array::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Arrayconst&)::{lambda(tvm::tir::Stmt const&)#1})
Array copy = arr.Map(fmutate);
^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 652, in tvm::runtime::Array::value, void>::type> tvm::runtime::Array::Map const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}) const
return Array(MapHelper(data_, fmap));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h", line 877, in tvm::runtime::ObjectPtr tvm::runtime::Array::MapHelper const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::runtime::ObjectPtr, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
U mapped = fmap(DowncastNoCheck(*it));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 210, in tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
auto fmutate = [self](const Stmt& s) { return self->VisitStmt(s); };
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 156, in tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
node.CopyOnWrite()->body = InjectOpaqueBlock(node->body, new_block_alloc_bufs);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc", line 219, in tvm::tir::BufferAllocationLocator::InjectOpaqueBlock(tvm::tir::Stmt, tvm::runtime::Array const&)
GetBlockReadWriteRegion(opaque_block, buffer_data_to_buffer_);
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 385, in tvm::tir::GetBlockReadWriteRegion(tvm::tir::Block const&, tvm::runtime::Map const&)
detector(block);
^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 135, in tvm::tir::BlockReadWriteDetector::operator()(tvm::tir::Stmt const&)
StmtExprVisitor::operator()(stmt);
^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 144, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BlockNode const*)
this->VisitStmt(op->body);
^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 119, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
VisitArray(op->seq, [this](const Stmt& s) { this->VisitStmt(s); });
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/functor_common.h", line 35, in VisitArray >
fvisit(arr[i]);
^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 119, in operator()
VisitArray(op->seq, [this](const Stmt& s) { this->VisitStmt(s); });
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 167, in tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::ForNode const*)
StmtVisitor::VisitStmt_(op);
^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 48, in tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::ForNode const*)
this->VisitStmt(op->body);
^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc", line 258, in tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::BlockRealizeNode const*)
Substitute(range->min, vmap), Substitute(range->extent, vmap))),
^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 758, in tvm::tir::Substitute(tvm::PrimExpr, std::function (tvm::tir::Var const&)>)
return IRSubstitute(vmap)(std::move(expr));
^^^^^^^^^^^^^^^^^^^^^^^^^^^
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 660, in tvm::tir::IRSubstitute::VisitExpr_(tvm::tir::VarNode const*)
ICHECK(ret_ex.dtype() == var.dtype()) << "substituting " << var << ":" << var.dtype()
^^^^^^^^^^^^^^^^^^^^^^^^^^^
tvm.error.InternalError: Traceback (most recent call last):
37: tvm::tir::transform::PrimFuncPassNode::operator()(tvm::IRModule, tvm::transform::PassContext const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/transform.cc:121
36: operator()
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:255
35: tvm::tir::PlanAndUpdateBufferAllocationLocation(tvm::tir::PrimFunc)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:247
34: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::BlockNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:178
33: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:211
32: tvm::runtime::Array::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:184
31: tvm::runtime::Array::value, void>::type> tvm::runtime::Array::Map const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}) const
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:652
30: tvm::runtime::ObjectPtr tvm::runtime::Array::MapHelper const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::runtime::ObjectPtr, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:877
29: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:210
28: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:145
27: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:211
26: tvm::runtime::Array::value, void>::type> tvm::tir::StmtMutator::Internal::MutateArray const&)::{lambda(tvm::tir::Stmt const&)#1}>(tvm::tir::StmtMutator*, tvm::runtime::Array::value, void>::type> const&, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:184
25: tvm::runtime::Array::value, void>::type> tvm::runtime::Array::Map const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}) const
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:652
24: tvm::runtime::ObjectPtr tvm::runtime::Array::MapHelper const&)::{lambda(tvm::tir::Stmt const&)#1}, tvm::tir::Stmt>(tvm::runtime::ObjectPtr, tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1})
at /data/qshenaf/envs/tvm/include/tvm/runtime/container/array.h:877
23: tvm::tir::StmtMutator::Internal::Mutate(tvm::tir::StmtMutator*, tvm::runtime::Array const&)::{lambda(tvm::tir::Stmt const&)#1}::operator()(tvm::tir::Stmt const&) const
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:210
22: tvm::tir::BufferAllocationLocator::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:156
21: tvm::tir::BufferAllocationLocator::InjectOpaqueBlock(tvm::tir::Stmt, tvm::runtime::Array const&)
at /data/qshenaf/envs/tvm/src/tir/transforms/plan_update_buffer_allocation_location.cc:219
20: tvm::tir::GetBlockReadWriteRegion(tvm::tir::Block const&, tvm::runtime::Map const&)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:385
19: tvm::tir::BlockReadWriteDetector::operator()(tvm::tir::Stmt const&)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:135
18: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::BlockNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:144
17: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::SeqStmtNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
16: VisitArray >
at /data/qshenaf/envs/tvm/src/tir/ir/functor_common.h:35
15: operator()
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:119
14: tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:167
13: tvm::tir::StmtVisitor::VisitStmt_(tvm::tir::ForNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:48
12: tvm::tir::BlockReadWriteDetector::VisitStmt_(tvm::tir::BlockRealizeNode const*)
at /data/qshenaf/envs/tvm/src/tir/analysis/block_access_region_detector.cc:258
11: tvm::tir::Substitute(tvm::PrimExpr, std::function (tvm::tir::Var const&)>)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:758
10: _ZThn16_N3tvm3tir15StmtExprMu
9: _ZThn16_N3tvm3tir15StmtExprMu
8: _ZThn16_N3tvm3tir15StmtExprMu
7: _ZThn16_N3tvm3tir15StmtExprMu
6: _ZThn16_N3tvm3tir15StmtExprMu
5: _ZThn16_N3tvm3tir15StmtExprMu
4: _ZThn16_N3tvm3tir15StmtExprMu
3: _ZThn16_N3tvm3tir15StmtExprMu
2: _ZThn16_N3tvm3tir15StmtExprMu
1: _ZThn16_N3tvm3tir12IRSubstitu
0: tvm::tir::IRSubstitute::VisitExpr_(tvm::tir::VarNode const*)
at /data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc:660
File "/data/qshenaf/envs/tvm/src/tir/ir/stmt_functor.cc", line 660
InternalError: Check failed: (ret_ex.dtype() == var.dtype()) is false: substituting v_ax4:int32 -> T.Ramp(0, 1, 3):int32x3

```

### Environment

tvm-0.21.dev0

### Steps to reproduce

```
import tvm
from tvm import te, topi, tir
from tvm import meta_schedule as ms

data = te.placeholder((2, 48, 28, 28, 3), dtype='float16', name='data')
gamma = te.placeholder((48,), dtype='float16', name='gamma')
beta = te.placeholder((48,), dtype='float16', name='beta')
op_config = {'data': data, 'gamma': gamma, 'beta': beta, 'num_groups': 12, 'channel_axis': 1, 'axes': [2, 3, 4], 'epsilon': 1e-08, }
op_output = topi.nn.group_norm(**op_config)

sch = tir.Schedule(te.create_prim_func([data, gamma, beta, op_output]).with_attr('target', tvm.target.Target('llvm')))

database = ms.tir_integration.tune_tir(sch.mod, target='llvm --num-cores=16', work_dir='./tune_tmp', max_trials_global=1, num_trials_per_iter=1)
sch = ms.tir_integration.compile_tir(database, sch.mod, 'llvm --num-cores=16')

passes = [tir.transform.VectorizeLoop(True),tir.transform.InjectVirtualThread()]
with tvm.ir.transform.PassContext(opt_level=4):
opt_mod = tvm.ir.transform.Sequential(passes)(sch.mod)
tvm.build(opt_mod, target='llvm')

```

### Triage

* needs-triage
* tir:transform
* tir:schedule
*tune:meta_schedule

Start with the reported topi.nn.group_norm_4.py reproduction and the tvm.build call, then trace the failure through src/tir/transforms/plan_update_buffer_allocation_location.cc and src/tir/analysis/block_access_region_detector.cc. Inspect the dtype substitution path at src/tir/ir/stmt_functor.cc and determine a focused regression test. Done means the reproduction builds without the reported InternalError and the regression test passes.

Written by the indexing model from the issue text.