Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

stablehlo.reshape op unsupported data type #1455

Open
mmanzoorTT opened this issue Dec 2, 2024 · 2 comments · May be fixed by #1626
Open

stablehlo.reshape op unsupported data type #1455

mmanzoorTT opened this issue Dec 2, 2024 · 2 comments · May be fixed by #1626
Assignees
@uazizTT
Labels
bug Something isn't working stablehlo conversion bug Bugs in StableHLO conversion
Milestone
[Third Party] HLO...

Comments

@mmanzoorTT
Copy link
Contributor

TTNN only supports reshape op for float and bfloat16. Pytorch models are generating transpose op for other data types.
@mmanzoorTT mmanzoorTT added this to the [Third Party] HLO + XLA milestone Dec 2, 2024
@mmanzoorTT mmanzoorTT added bug Something isn't working stablehlo conversion bug Bugs in StableHLO conversion labels Dec 2, 2024
@mmanzoorTT
Copy link
Contributor Author

TTIR graphs

#any_device_tile = #tt.operand_constraint<dram|l1|tile|none|interleaved|single_bank|height_sharded|width_sharded|block_sharded|any_layout|any_device_tile>
module {
  func.func @main(%arg0: tensor<1x1x12x16xi32>) -> tensor<1x192xi32> {
    %0 = tensor.empty() : tensor<1x192xi32>
    %1 = "ttir.reshape"(%arg0, %0) <{operand_constraints = [#any_device_tile, #any_device_tile], shape = [1 : i32, 192 : i32]}> : (tensor<1x1x12x16xi32>, tensor<1x192xi32>) -> tensor<1x192xi32>
    return %1 : tensor<1x192xi32>
  }
}

@mmanzoorTT
Copy link
Contributor Author

TTRT dump

�[32m                 Always�[0m | �[1m�[38;5;208m WARNING�[0m | Golden information not found
�[38;2;000;128;000m                 Always�[0m | �[1m�[38;2;255;000;000mFATAL   �[0m | Error
2024-11-29 01:56:14,687 - INFO - set log file=
2024-11-29 01:56:14,687 - INFO - setting artifacts folder path=/localdev/mmanzoor/compiler/tt-torch/ttrt-artifacts
2024-11-29 01:56:14,687 - DEBUG - ----------------------------starting run API----------------------------
2024-11-29 01:56:14,687 - DEBUG - ------preprocessing run API
2024-11-29 01:56:14,687 - DEBUG - ----------------------------starting query API----------------------------
2024-11-29 01:56:14,687 - DEBUG - ------preprocessing query API
2024-11-29 01:56:14,687 - DEBUG - ------finished preprocessing query API
2024-11-29 01:56:14,687 - DEBUG - ------checking constraints for query API
2024-11-29 01:56:14,687 - DEBUG - ------finished constraints for query API
2024-11-29 01:56:14,687 - DEBUG - ------executing query API
2024-11-29 01:56:14,693 - DEBUG - getting system descriptor
2024-11-29 01:56:16,281 - DEBUG - ------finished executing query API
2024-11-29 01:56:16,281 - DEBUG - ------postprocessing query API
2024-11-29 01:56:16,281 - INFO - PASS: getting system_desc passed
2024-11-29 01:56:16,282 - INFO - results saved to=query_results.json
2024-11-29 01:56:16,283 - DEBUG - ------finished postprocessing query API
2024-11-29 01:56:16,283 - DEBUG - ----------------------------finished query API----------------------------
2024-11-29 01:56:16,283 - DEBUG - ------finished preprocessing read API
2024-11-29 01:56:16,283 - DEBUG - ------checking constraints for run API
2024-11-29 01:56:16,283 - DEBUG - finding all ttnn files from=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - checking if file=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn exists
2024-11-29 01:56:16,283 - DEBUG - getting file extension of=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - found file=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - finding all ttmetal files from=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - checking if file=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn exists
2024-11-29 01:56:16,283 - DEBUG - getting file extension of=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - ttnn_binary_paths=['/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn']
2024-11-29 01:56:16,283 - DEBUG - ttmetal_binary_paths=[]
2024-11-29 01:56:16,283 - DEBUG - getting file name of=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,283 - DEBUG - getting file extension of=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,835 - DEBUG - ------finished checking constraints for run API
2024-11-29 01:56:16,835 - DEBUG - ------executing run API
2024-11-29 01:56:16,835 - DEBUG - executing ttnn binaries
2024-11-29 01:56:16,835 - DEBUG - setting tt runtime debug env=debug::Env{
	loadKernelsFromDisk: 0,
	enableAsyncTTNN: 0
}
2024-11-29 01:56:16,835 - DEBUG - setting tt runtime workaround env=workaround::Env{
	ignoreTileShape: 1,
	emptyOpForceRowMajor: 1,
	fullOpForceRowMajor: 1,
	maxpool2dPreshard: 1,
	swapBinaryOperands: 1
}
2024-11-29 01:56:16,835 - DEBUG - setting torch manual seed=0
2024-11-29 01:56:16,836 - DEBUG - opening devices=[0, 1]
2024-11-29 01:56:16,876 - INFO - evaluating binary=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,876 - DEBUG - evaluating program=0 for binary=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,877 - DEBUG - generating inputs/outputs for loop=1/1 for binary=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:16,877 - DEBUG - starting loop=1/1 for binary=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:17,062 - ERROR - ERROR: test=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn experienced an error with exception=TT_FATAL @ /localdev/mmanzoor/compiler/tt-torch/third_party/tt-mlir/src/tt-mlir/third_party/tt-metal/src/tt-metal/ttnn/cpp/ttnn/operations/data_movement/reshape_on_device/device/reshape_op.cpp:21: input_tensor_a.get_dtype() == DataType::BFLOAT16 or input_tensor_a.get_dtype() == DataType::FLOAT32
info:
Error
backtrace:
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/lib/python3.11/site-packages/ttrt/runtime/_C.cpython-311-x86_64-linux-gnu.so(+0x94f98) [0x7ff83fcd3f98]
 --- ttnn::operations::data_movement::ReshapeDeviceOperation::validate(std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > const&) const
 --- void ttnn::device_operation::detail::launch_on_worker_thread<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >, unsigned char, long, tt::tt_metal::operation::DeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_args_t, std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> >, tt::tt_metal::v0::Device*>(unsigned char, long, tt::tt_metal::operation::DeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > > const&, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_args_t const&, std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> >&, tt::tt_metal::v0::Device*&)
 --- tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_return_value_t ttnn::device_operation::detail::launch_on_single_device<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > > >(unsigned char, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::operation_attributes_t const&, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_args_t const&)
 --- /localdev/mmanzoor/compiler/tt-torch/install/lib/_ttnn.so(+0x1dd35c9) [0x7ff83f48f5c9]
 --- tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_return_value_t ttnn::device_operation::detail::invoke<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > > >(unsigned char, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::operation_attributes_t const&, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >::tensor_args_t const&)
 --- /localdev/mmanzoor/compiler/tt-torch/install/lib/_ttnn.so(+0x1dcf79f) [0x7ff83f48b79f]
 --- std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > tt::tt_metal::operation::run<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >(tt::tt_metal::operation::DeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > >&&, std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor> > const&, std::vector<std::optional<tt::tt_metal::Tensor const>, std::allocator<std::optional<tt::tt_metal::Tensor const> > > const&, std::vector<std::optional<tt::tt_metal::Tensor>, std::allocator<std::optional<tt::tt_metal::Tensor> > > const&, unsigned char)
 --- ttnn::operations::data_movement::ReshapeOperation::invoke(unsigned char, tt::tt_metal::Tensor const&, ttnn::types::Shape, std::optional<tt::tt_metal::MemoryConfig> const&)
 --- /localdev/mmanzoor/compiler/tt-torch/install/lib/_ttnn.so(+0x1c5a01e) [0x7ff83f31601e]
 --- /localdev/mmanzoor/compiler/tt-torch/install/lib/_ttnn.so(+0x1c30c89) [0x7ff83f2ecc89]
 --- ttnn::operations::data_movement::detail::convert_tensor_to_rm_reshape_convert_back_to_orig_layout(tt::tt_metal::Tensor const&, ttnn::types::Shape const&)
 --- ttnn::operations::data_movement::ReshapeViewOperation::invoke(tt::tt_metal::Tensor const&, ttnn::types::Shape const&)
 --- ttnn::operations::data_movement::ReshapeViewOperation::invoke(tt::tt_metal::Tensor const&, tt::stl::Span<int const, 18446744073709551615ul>)
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/lib/python3.11/site-packages/ttrt/runtime/_C.cpython-311-x86_64-linux-gnu.so(+0xd890b) [0x7ff83fd1790b]
 --- tt::runtime::ttnn::operations::data_movement::run(tt::target::ttnn::ReshapeOp const*, tt::runtime::ttnn::ProgramContext&)
 --- tt::runtime::ttnn::runProgram(tt::tt_metal::distributed::MeshDevice&, tt::runtime::Binary&, unsigned int, std::vector<tt::tt_metal::Tensor*, std::allocator<tt::tt_metal::Tensor*> > const&, std::vector<tt::tt_metal::Tensor*, std::allocator<tt::tt_metal::Tensor*> > const&)
 --- tt::runtime::ttnn::submit(tt::runtime::Device, tt::runtime::Binary, unsigned int, std::vector<tt::runtime::Tensor, std::allocator<tt::runtime::Tensor> > const&, std::vector<tt::runtime::Tensor, std::allocator<tt::runtime::Tensor> > const&)
 --- tt::runtime::submit(tt::runtime::Device, tt::runtime::Binary, unsigned int, std::vector<tt::runtime::Tensor, std::allocator<tt::runtime::Tensor> > const&, std::vector<tt::runtime::Tensor, std::allocator<tt::runtime::Tensor> > const&)
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/lib/python3.11/site-packages/ttrt/runtime/_C.cpython-311-x86_64-linux-gnu.so(+0x5a175) [0x7ff83fc99175]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/lib/python3.11/site-packages/ttrt/runtime/_C.cpython-311-x86_64-linux-gnu.so(+0x7c09b) [0x7ff83fcbb09b]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/lib/python3.11/site-packages/ttrt/runtime/_C.cpython-311-x86_64-linux-gnu.so(+0x6fb27) [0x7ff83fcaeb27]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x51ad67]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyObject_MakeTpCall+0x22c) [0x4e75dc]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyEval_EvalFrameDefault+0x8f2) [0x4fb152]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyObject_FastCallDictTstate+0x1b1) [0x4eefb1]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyObject_Call_Prepend+0x59) [0x53c709]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x5cbea7]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyObject_MakeTpCall+0x22c) [0x4e75dc]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyEval_EvalFrameDefault+0x8f2) [0x4fb152]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x62e1b4]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(PyEval_EvalCode+0x97) [0x4f3a67]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x647c37]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x645350]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python() [0x650d15]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyRun_SimpleFileObject+0x194) [0x650a64]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_PyRun_AnyFileObject+0x43) [0x650833]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(Py_RunMain+0x327) [0x64f787]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(Py_BytesMain+0x2d) [0x61ee0d]
 --- /lib/x86_64-linux-gnu/libc.so.6(+0x29d90) [0x7ff840f75d90]
 --- /lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x80) [0x7ff840f75e40]
 --- /localdev/mmanzoor/compiler/tt-torch/env/venv/bin/python(_start+0x25) [0x61ec95]

2024-11-29 01:56:17,064 - DEBUG - finished executing ttnn binaries
2024-11-29 01:56:17,064 - DEBUG - executing ttmetal binaries
2024-11-29 01:56:17,064 - WARNING - no binaries found to run - returning early
2024-11-29 01:56:17,064 - DEBUG - finished executing ttmetal binaries
2024-11-29 01:56:17,064 - DEBUG - ------finished executing run API
2024-11-29 01:56:17,064 - DEBUG - ------postprocessing run API
2024-11-29 01:56:17,064 - ERROR - ERROR: test case=/localdev/mmanzoor/compiler/tt-torch/results/mlir_tests/output/ttir/aten::view_4843.ttnn
2024-11-29 01:56:17,065 - INFO - results saved to=run_results.json
2024-11-29 01:56:17,065 - DEBUG - ------finished postprocessing run API
2024-11-29 01:56:17,065 - DEBUG - ----------------------------finished run API----------------------------

@uazizTT uazizTT self-assigned this Dec 12, 2024
@uazizTT uazizTT linked a pull request Dec 18, 2024 that will close this issue
Draft
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@uazizTT uazizTT

Labels
bug Something isn't working stablehlo conversion bug Bugs in StableHLO conversion
Projects
None yet
Milestone
[Third Party] HLO + XLA
Development

Successfully merging a pull request may close this issue.

Introduce a workaround to add typecast for unsupported types tenstorrent/tt-mlir
2 participants
@mmanzoorTT @uazizTT