Convolution lowering causes runtime crash in multiple models

[nemanjagrujic]

Convolution lowering causes runtime crash in multiple models:

• retinanet_resnet50_fpn
• retinanet_resnet50_fpn_v2
• ssd300_vgg16
• vgg19_bn

This method:

def insert_sharded_nxc_to_ncx(g, output_tensor, output_shape):
    output_tensor = g.call_function(ttnn.sharded_to_interleaved, (output_tensor, TtnnL1MemoryConfig()))
    target_shape = (output_shape[0], *output_shape[2:], output_shape[1])
    output_tensor = g.call_function(ttnn.reshape, (output_tensor, target_shape))
    target_permute = [0, len(output_shape) - 1] + list(range(1, len(output_shape) - 1))
    return g.call_function(ttnn.permute, (output_tensor, target_permute))

Puts tensor in L1 which later causes runtime error for large tensors:

RuntimeError: TT_THROW @ /home/ubuntu/tt-metal/tt_metal/impl/allocator/bank_manager.cpp:132: tt::exception
info:
Out of Memory: Not enough space to allocate 8257536 B L1 buffer across 64 banks, where each bank needs to store 129024 B
backtrace:
 --- /home/ubuntu/tt-metal/build_Release/lib/libtt_metal.so(+0x205c85) [0x7fadd3c53c85]
 --- tt::tt_metal::BankManager::allocate_buffer(unsigned long, unsigned long, bool, CoreRangeSet const&, std::__1::optional<unsigned int>)
 --- tt::tt_metal::Allocator::allocate_buffer(tt::tt_metal::Buffer*)
 --- tt::tt_metal::Buffer::allocate_impl()
 --- /home/ubuntu/tt-metal/build_Release/lib/libtt_metal.so(+0x1a1670) [0x7fadd3bef670]
 --- /home/ubuntu/tt-metal/build_Release/lib/libtt_metal.so(+0x187a71) [0x7fadd3bd5a71]
 --- tt::tt_metal::Buffer::create(tt::tt_metal::IDevice*, unsigned long, unsigned long, tt::tt_metal::BufferType, tt::tt_metal::TensorMemoryLayout, std::__1::optional<tt::tt_metal::ShardSpecBuffer> const&, std::__1::optional<bool>, std::__1::optional<tt::stl::StrongType<unsigned char, tt::tt_metal::SubDeviceIdTag>>)
 --- tt::tt_metal::tensor_impl::allocate_buffer_on_device(tt::tt_metal::IDevice*, tt::tt_metal::TensorSpec const&)
 --- tt::tt_metal::create_device_tensor(tt::tt_metal::TensorSpec const&, tt::tt_metal::IDevice*)
 --- ttnn::operations::data_movement::PermuteDeviceOperation::create_output_tensors(ttnn::operations::data_movement::PermuteDeviceOperation::operation_attributes_t const&, ttnn::operations::data_movement::PermuteDeviceOperation::tensor_args_t const&)
 --- ttnn::operations::data_movement::PermuteDeviceOperation::tensor_return_value_t ttnn::device_operation::detail::launch_on_single_device<ttnn::operations::data_movement::PermuteDeviceOperation>(tt::stl::StrongType<unsigned char, ttnn::QueueIdTag>, ttnn::operations::data_movement::PermuteDeviceOperation::operation_attributes_t const&, ttnn::operations::data_movement::PermuteDeviceOperation::tensor_args_t const&)
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x1b5feba) [0x7fada6935eba]
 --- ttnn::operations::data_movement::PermuteDeviceOperation::tensor_return_value_t ttnn::device_operation::detail::invoke<ttnn::operations::data_movement::PermuteDeviceOperation>(tt::stl::StrongType<unsigned char, ttnn::QueueIdTag>, ttnn::operations::data_movement::PermuteDeviceOperation::operation_attributes_t const&, ttnn::operations::data_movement::PermuteDeviceOperation::tensor_args_t const&)
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x1bd5ba9) [0x7fada69abba9]
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x1bd57c8) [0x7fada69ab7c8]
 --- ttnn::operations::data_movement::detail::permute_impl(tt::tt_metal::Tensor const&, tt::stl::SmallVector<unsigned int, 8ul> const&, tt::tt_metal::MemoryConfig const&, std::__1::optional<float> const&)
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x1bd8ec1) [0x7fada69aeec1]
 --- void tt::tt_metal::operation::launch_op_func<std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>>(std::__1::function<std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> (std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>> const&)> const&, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>>, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>>)
 --- void tt::tt_metal::operation::launch_op<std::__1::function<std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> (std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>> const&)>, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>>(std::__1::function<std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> (std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>> const&)>&&, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>>, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>>)
 --- tt::tt_metal::operation::launch_with_autoformat(std::__1::function<std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> (std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>> const&)>&&, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>> const&, std::__1::vector<tt::tt_metal::Tensor, std::__1::allocator<tt::tt_metal::Tensor>>&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor const>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor const>>> const&, std::__1::vector<std::__1::optional<tt::tt_metal::Tensor>, std::__1::allocator<std::__1::optional<tt::tt_metal::Tensor>>> const&)
 --- ttnn::operations::data_movement::detail::permute_launch(tt::tt_metal::Tensor const&, tt::stl::SmallVector<unsigned int, 8ul> const&, tt::tt_metal::MemoryConfig const&, std::__1::optional<float> const&)
 --- ttnn::operations::data_movement::ExecutePermute::invoke(tt::stl::StrongType<unsigned char, ttnn::QueueIdTag>, tt::tt_metal::Tensor const&, tt::stl::SmallVector<long, 8ul> const&, std::__1::optional<tt::tt_metal::MemoryConfig> const&, std::__1::optional<float> const&)
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x199ee7b) [0x7fada6774e7b]
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x19950be) [0x7fada676b0be]
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0x1994efa) [0x7fada676aefa]
 --- /home/ubuntu/tt-metal/ttnn/ttnn/_ttnn.so(+0xebf59e) [0x7fada5c9559e]
 --- python(PyCFunction_Call+0x59) [0x5f3439]
 --- python(_PyObject_MakeTpCall+0x29e) [0x5f38ce]
 --- python() [0x50ac38]
 --- python(PyObject_Call+0x1f7) [0x5f2ff7]
 --- python() [0x59cb2c]
 --- python(PyObject_Call+0x27e) [0x5f307e]
 --- python(_PyEval_EvalFrameDefault+0x1f44) [0x56c764]
 --- python(_PyEval_EvalCodeWithName+0x26a) [0x56921a]
 --- python(_PyFunction_Vectorcall+0x393) [0x5f6893]
 --- python() [0x59c79e]
 --- python(_PyObject_MakeTpCall+0x29e) [0x5f38ce]
 --- python(_PyEval_EvalFrameDefault+0x5de2) [0x570602]
 --- python(_PyFunction_Vectorcall+0x1b6) [0x5f66b6]
 --- python(PyObject_Call+0x1f7) [0x5f2ff7]
 --- python(_PyEval_EvalFrameDefault+0x1f44) [0x56c764]
 --- python(_PyEval_EvalCodeWithName+0x26a) [0x56921a]
 --- python(PyEval_EvalCode+0x27) [0x68ce37]
 --- python() [0x67c4e1]
 --- python() [0x67c55f]
 --- python() [0x67c601]
 --- python(PyRun_SimpleFileExFlags+0x197) [0x67e727]
 --- python(Py_RunMain+0x212) [0x6b6e72]
 --- python(Py_BytesMain+0x2d) [0x6b71fd]
 --- /lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0xf3) [0x7fade4931083]
 --- python(_start+0x2e) [0x5fa55e]

• Steps to reproduce:

1. Generate code for torch vision object detection models: pytest tests/models/torchvision/test_torchvision_object_detection.py --gen_op_accuracy_tests
2. Run the models: python tests/autogen_accuracy_tests/ssd300_vgg16_code.py or tests/autogen_accuracy_tests/retinanet_resnet50_fpn_code.py or tests/autogen_accuracy_tests/retinanet_resnet50_fpn_v2_code.py

[nemanjagrujic]

@ayerofieiev-tt @kevinwuTT This problem crashes multiple models. There should be some check if tensor is too large for L1. Also consider to introduce some strategy to get the tensor back from L1 to DRAM if its to large for operation (eg operation makes tensor even larger)