You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by "Cookiee235 (via GitHub)" <gi...@apache.org> on 2024/08/30 13:17:17 UTC
[I] [Bug] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure [tvm]
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: commits-unsubscribe@tvm.apache.org.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org
Re: [I] [Bug] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure [tvm]
Posted by "Cookiee235 (via GitHub)" <gi...@apache.org>.
Cookiee235 commented on issue #17322:
URL: https://github.com/apache/tvm/issues/17322#issuecomment-2329531089
@Lunderberg Thanks for your deep investigation.
This error is caused by an invalid input value (i.e., 1025), and it cannot be raised at compile time. Just moving the error to an earlier point during runtime is not very important. Thus, I will close this 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] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure [tvm]
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on issue #17322:
URL: https://github.com/apache/tvm/issues/17322#issuecomment-2326984086
That said, it may still be possible to raise an error at an earlier point during runtime. If this error occurred at the end of a very long end-to-end model execution, then it may be beneficial to move the implicit size validation of `R.builtin.alloc_tensor` into an explicit size validation at the start of the function. I could see this either being based on some hypothetical `CollectImplicitAssumptions` utility, which would identify `offset + requested_size <= backing_buffer_size` as a requirement, or by normalization of an operator to make new assert statements if their preconditions cannot already be proven.
--
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] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure [tvm]
Posted by "Lunderberg (via GitHub)" <gi...@apache.org>.
Lunderberg commented on issue #17322:
URL: https://github.com/apache/tvm/issues/17322#issuecomment-2326946938
Reading through the function, the out-of-bounds access would occur any time the input array is larger than 16, and is caught during the call to `R.memory.alloc_storage`. If this were written using `R.call_tir`, then the dynamic shape of the output would be used to generate a dynamic size for `R.memory.alloc_storage`. However, the static size in `R.memory.alloc_storage` implicitly assumes that `m <= 16`.
(Also, the error first occurs at 1025 rather than 17, because 1025 is the first time that the 4kB page size defined by `PooledAllocated::kDefaultPageSize` would be exceeded. The `R.memory.alloc_storage` may return a larger size than requested, in order to re-use similarly sized allocations. The validation done by `R.memory.alloc_tensor` is then based on the actual size rather than the requested size.)
> If the given Relax IR is illegal, do we have a corresponding legality-checking mechanism to alert early?
I think the Relax IR is legal, even though the error can be triggered. In general, the Relax validation tries to categorize IR as either (a) provably incorrect for all possible values, (b) provably correct for all possible values, or (c) neither of the above. If something falls into category (a), then a compile-time error can be thrown. If something falls into category (b), then a runtime check may be omitted. If something falls into (c), then we need to wait until runtime to perform validation.
In this case, because the runtime value of `m` may be 16 or less, we cannot raise a compile-time error for this IR, and catching the error at runtime is the best option.
--
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] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure [tvm]
Posted by "Cookiee235 (via GitHub)" <gi...@apache.org>.
Cookiee235 closed issue #17322: [Bug] InternalError: Check failed: (offset + needed_size <= this->buffer.size) is false: storage allocation failure
URL: https://github.com/apache/tvm/issues/17322
--
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