Cookiee235 opened a new issue, #17322: URL: https://github.com/apache/tvm/issues/17322
The relax IR crashed when the input tensor size was larger than 1024. If this `tir` is illegal, do we have a corresponding legality-checking mechanism to alert early? A similar issue is found here: https://github.com/mlc-ai/mlc-llm/issues/971 ### Actual behavior ``` Traceback (most recent call last): File "/share_container/optfuzz/res/res_0830/optfuzz_nnsmith/res_executions/881_test.py", line 260, in <module> before_outputs, infer_time1 = compile_mod(mod, 'main', 'llvm', input_0,input_1,input_2) ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/share_container/optfuzz/res/res_0830/optfuzz_nnsmith/res_executions/881_test.py", line 251, in compile_mod mod_outputs = vm[f'{func_name}'](*inputs) ^^^^^^^^^^^^^^^^^^^^^^^^^^^ File "/software/tvm/python/tvm/_ffi/_ctypes/packed_func.py", line 239, in __call__ raise_last_ffi_error() File "/software/tvm/python/tvm/_ffi/base.py", line 481, in raise_last_ffi_error raise py_err tvm.error.InternalError: Traceback (most recent call last): 14: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::_LookupFunction(tvm::runtime::String const&)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 13: tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 12: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 11: tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) 10: tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() 9: tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) 8: tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 7: tvm::runtime::PackedFuncObj::Extractor<tvm::runtime::PackedFuncSubObj<tvm::runtime::relax_vm::VirtualMachineImpl::GetClosureInternal(tvm::runtime::String const&, bool)::{lambda(tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*)#1}> >::Call(tvm::runtime::PackedFuncObj const*, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 6: tvm::runtime::relax_vm::VirtualMachineImpl::InvokeBytecode(long, std::vector<tvm::runtime::TVMRetValue, std::allocator<tvm::runtime::TVMRetValue> > const&) 5: tvm::runtime::relax_vm::VirtualMachineImpl::RunLoop() 4: tvm::runtime::relax_vm::VirtualMachineImpl::RunInstrCall(tvm::runtime::relax_vm::VMFrame*, tvm::runtime::relax_vm::Instruction) 3: tvm::runtime::relax_vm::VirtualMachineImpl::InvokeClosurePacked(tvm::runtime::ObjectRef const&, tvm::runtime::TVMArgs, tvm::runtime::TVMRetValue*) 2: _ZN3tvm7runtime13PackedFuncObj9ExtractorINS0_1 1: tvm::runtime::TypedPackedFunc<tvm::runtime::NDArray (tvm::runtime::memory::Storage, long, tvm::runtime::ShapeTuple, DLDataType)>::AssignTypedLambda<tvm::runtime::Registry::set_body_method<tvm::runtime::memory::Storage, tvm::runtime::memory::StorageObj, tvm::runtime::NDArray, long, tvm::runtime::ShapeTuple, DLDataType, void>(tvm::runtime::NDArray (tvm::runtime::memory::StorageObj::*)(long, tvm::runtime::ShapeTuple, DLDataType))::{lambda(tvm::runtime::memory::Storage, long, tvm::runtime::ShapeTuple, DLDataType)#1}>(tvm::runtime::Registry::set_body_method<tvm::runtime::memory::Storage, tvm::runtime::memory::StorageObj, tvm::runtime::NDArray, long, tvm::runtime::ShapeTuple, DLDataType, void>(tvm::runtime::NDArray (tvm::runtime::memory::StorageObj::*)(long, tvm::runtime::ShapeTuple, DLDataType))::{lambda(tvm::runtime::memory::Storage, long, tvm::runtime::ShapeTuple, DLDataType)#1}, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >)::{lambda(tvm::runti me::TVMArgs const&, tvm::runtime::TVMRetValue*)#1}::operator()(tvm::runtime::TVMArgs const, tvm::runtime::TVMRetValue) const 0: tvm::runtime::memory::StorageObj::AllocNDArray(long, tvm::runtime::ShapeTuple, DLDataType) File "/software/tvm/src/runtime/memory/memory_manager.cc", line 117 InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure, attempted to allocate 36480 at offset 0 in region that is 4096bytes ``` ### Steps to reproduce ``` import tvm from tvm import relax import numpy as np from tvm.script import ir as I from tvm.script import tir as T from tvm.script import relax as R @I.ir_module class Module: @T.prim_func def add_one(x_handle: T.handle, y_handle: T.handle): m = T.int64() x = T.match_buffer(x_handle, (m,)) y = T.match_buffer(y_handle, (m,)) # with T.block("root"): for i in range(m): with T.block("add"): vi = T.axis.spatial(m, i) T.reads(x[vi]) T.writes(y[vi]) y[vi] = x[vi] + T.float32(1) @R.function def main(x: R.Tensor(("m",), dtype="float32")) -> R.Tensor(("m",), dtype="float32"): m = T.int64() R.func_attr({"relax.force_pure": 1, "relax.rewrite_cuda_graph.capture_symbolic_vars": ["m"]}) cls = Module storage: R.Object = R.memory.alloc_storage(R.shape([16]), R.prim_value(0), R.str("global"), R.dtype("float32")) alloc1: R.Tensor((m,), dtype="float32") = R.memory.alloc_tensor(storage, R.prim_value(0), R.shape([m]), R.dtype("float32")) cls.add_one(x, alloc1) storage1: R.Object = R.memory.alloc_storage(R.shape([16]), R.prim_value(0), R.str("global"), R.dtype("float32")) alloc2: R.Tensor((m,), dtype="float32") = R.memory.alloc_tensor(storage1, R.prim_value(0), R.shape([m]), R.dtype("float32")) cls.add_one(alloc1, alloc2) alloc3: R.Tensor((m,), dtype="float32") = R.builtin.alloc_tensor(R.shape([m]), R.dtype("float32"), R.prim_value(0), R.str("global")) cls.add_one(alloc2, alloc3) return alloc3 mod = Module mod = tvm.tir.analysis.OOBChecker()(mod) def compile_mod(mod, func_name, target, *inputs): ex = relax.build(mod, target='llvm') vm = relax.VirtualMachine(ex, tvm.cpu()) mod_outputs = vm[f'{func_name}'](*inputs) mod_outputs = mod_outputs.numpy() input_ = tvm.nd.array(np.random.random([1025]).astype('float32')) # tensor_size > 1024 will lead to crash compile_mod(mod, 'main', 'llvm', input_) ``` CC @Lunderberg @junrushao -- 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: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
