search

tenstorrent / tt-metal

:metal: TT-NN operator library, and TT-Metalium low level kernel programming model.
Apache License 2.0
475 stars 75 forks source link

[Bug Report] Wrong fields in tracy report #15209

Open skrsticTT opened 14 hours ago

skrsticTT commented 14 hours ago

By running this simple test in python:

import pytest
import torch
import ttnn

@pytest.mark.parametrize("dims", [(32, 32), (64, 64)])
def test_add_with_block_sharding(device, dims):
    torch.manual_seed(0)
    h = dims[0]
    w = dims[1]
    torch_input_tensor_a = torch.rand((h, w))
    torch_input_tensor_b = torch.rand((h, w))

    input_tensor_a = ttnn.from_torch(
        torch_input_tensor_a,
        memory_config=ttnn.DRAM_MEMORY_CONFIG,
        layout=ttnn.TILE_LAYOUT,
        device=device,
        dtype=ttnn.bfloat16
    )

    input_tensor_b = ttnn.from_torch(
        torch_input_tensor_b,
        memory_config=ttnn.DRAM_MEMORY_CONFIG,
        layout=ttnn.TILE_LAYOUT,
        device=device,
        dtype=ttnn.bfloat16,
    )

    output = ttnn.add(input_tensor_a, input_tensor_b, memory_config=ttnn.DRAM_MEMORY_CONFIG)

I got wrong input dims for the second test:

INPUT_0_Y | INPUT_0_X
-- | --
32 | 32
32 | 32

It also affects second input and output.

Not that all of 3 tensors are interleaved. If I change output to be sharded for example, everything is ok.

mo-tenstorrent commented 12 hours ago

Can you provide the content of the report folder generated after you create the csv report.

It would be under generated/profiler/report/

skrsticTT commented 12 hours ago

tracy_data.zip

Here is data from report folder.

mo-tenstorrent commented 11 hours ago

This is a cache hit issue, for some reason both ops are giving the same hash.

As a quick remedy, running the test with python -m tracy -r -v -p --no-op-info-cache -m pytest {test_path} Will fix the above issue.

Note that in real tests, disabling op info cache will add considerable overhead host side time for the op.

mo-tenstorrent commented 11 hours ago

Looking into why both ops are giving the same hash.

skrsticTT commented 10 hours ago

Thanks Mo! Btw where I can check for an op hash?

mo-tenstorrent commented 9 hours ago

The perf report doesn't report it, it is buried in op infra C++ code.

mo-tenstorrent commented 7 hours ago

As per conversation with @dmakoviichuk-tt and @yan-zaretskiy There are ops that ignore IO tensor shape in their compute has function. While ideal for program cache, this is not good for op info caching as it causes issues observed in this issue.

mo-tenstorrent commented 7 hours ago

One idea is to use the default hash function here

First try, it is running into build issues:

FAILED: ttnn/CMakeFiles/ttnn.dir/Unity/unity_0_cxx.cxx.o
/usr/bin/clang++-17 -DARCH_WORMHOLE_B0 -DBOOST_CONTAINER_NO_LIB -DBOOST_CONTAINER_STATIC_LINK -DFMT_HEADER_ONLY=1 -DTRACY_ENABLE -DTRACY_IMPORTS -DTTNN_WITH_PYTHON_BINDINGS=1 -DUSING_UV_SHARED=1 -Dttnn_EXPORTS -I/home/mmemarian/infra/tt-metal/tt_metal/third_party/umd -I/home/mmemarian/infra/tt-metal -I/home/mmemarian/infra/tt-metal/tt_metal -I/home/mmemarian/infra/tt-metal/ttnn -I/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/deprecated -I/home/mmemarian/infra/tt-metal/ttnn/cpp -I/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/experimental/tt_lib -I/home/mmemarian/infra/tt-metal/tt_metal/third_party/pybind11/include -I/usr/include/python3.8 -I/home/mmemarian/infra/tt-metal/tt_metal/include -I/home/mmemarian/infra/tt-metal/tt_metal/impl -I/home/mmemarian/infra/tt-metal/tt_metal/tt_stl -I/home/mmemarian/infra/tt-metal/tt_metal/third_party/umd/device -I/home/mmemarian/infra/tt-metal/.cpmcache/nanomsg/28cc32d5bdb6a858fe53b3ccf7e923957e53eada/include -I/home/mmemarian/infra/tt-metal/.cpmcache/flatbuffers/2c4062bffa52fa4157b1b4deeae73395df475fda/include -I/home/mmemarian/infra/tt-metal/.cpmcache/libuv/72aa6fc919b0653b47b7e9e7cf59ec92037e8693/include -I/home/mmemarian/infra/tt-metal/tt_metal/hw/inc -I/home/mmemarian/infra/tt-metal/tt_metal/hw/inc/wormhole -I/home/mmemarian/infra/tt-metal/tt_metal/hw/inc/wormhole/wormhole_b0_defines -I/home/mmemarian/infra/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/wormhole -isystem /home/mmemarian/infra/tt-metal/.cpmcache/magic_enum/1e1af177d4ab0ef660f105434fd1017c4d1f8c17/include/magic_enum -isystem /home/mmemarian/infra/tt-metal/.cpmcache/fmt/73b5ec45edbd92babfd91c3777a9e1ab9cac8238/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_core/e679bef5c160cf29d0f37d549881dc5f5a58c332/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_assert/3ab1f6f9db9a884ad9a641164dbb6589a5aa7e2d/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_config/0bad5ba3b48288a243894aa801ed6eccbef70b60/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_static_assert/4be2778cf9cc81907a60618aede8cc2794ead5e2/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_throw_exception/af21673c57a398b5939be5615af8d2614136f98b/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f -isystem /home/mmemarian/infra/tt-metal/tt_metal/third_party/tracy/public -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_container/5fb02b14b46d0d84e7a0ce09e2ea5e963d5d93bd/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_intrusive/4a7bf962355d8580809cea3c68f55bbaaa746e64/include -isystem /home/mmemarian/infra/tt-metal/.cpmcache/boost_move/c59effd88face3140123440bc5425ee60328f08d/include -O3 -std=c++20 -fPIC -fvisibility=default -stdlib=libc++ -Werror -Wdelete-non-virtual-dtor -Wreturn-type -Wswitch -Wuninitialized -Wno-unused-parameter -mavx2 -fPIC -fvisibility-inlines-hidden -fno-lto -Wsometimes-uninitialized -Wno-c++11-narrowing -Wno-error=local-type-template-args -Wno-delete-non-abstract-non-virtual-dtor -Wno-c99-designator -Wno-shift-op-parentheses -Wno-non-c-typedef-for-linkage -Wno-deprecated-this-capture -Wno-deprecated-volatile -Wno-deprecated-builtins -Wno-deprecated-declarations -fno-omit-frame-pointer -MP -Wno-int-to-pointer-cast -fno-var-tracking -Winvalid-pch -Xclang -include-pch -Xclang /home/mmemarian/infra/tt-metal/build_Release_tracy/ttnn/CMakeFiles/ttnn.dir/cmake_pch.hxx.pch -Xclang -include -Xclang /home/mmemarian/infra/tt-metal/build_Release_tracy/ttnn/CMakeFiles/ttnn.dir/cmake_pch.hxx -MD -MT ttnn/CMakeFiles/ttnn.dir/Unity/unity_0_cxx.cxx.o -MF ttnn/CMakeFiles/ttnn.dir/Unity/unity_0_cxx.cxx.o.d -o ttnn/CMakeFiles/ttnn.dir/Unity/unity_0_cxx.cxx.o -c /home/mmemarian/infra/tt-metal/build_Release_tracy/ttnn/CMakeFiles/ttnn.dir/Unity/unity_0_cxx.cxx
In file included from <built-in>:1:
In file included from /home/mmemarian/infra/tt-metal/build_Release_tracy/ttnn/CMakeFiles/ttnn.dir/cmake_pch.hxx:5:
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1124:13: error: static assertion failed due to requirement 'tt::stl::concepts::always_false_v<tt::tt_metal::operation::DeviceOperation<std::vector<tt::tt_metal::Tensor, std::allocator<tt::tt_metal::Tensor>>>>': Type doesn't support hashing using tt::stl::hash::hash_object
 1124 |             tt::stl::concepts::always_false_v<T>, "Type doesn't support hashing using tt::stl::hash::hash_object");
      |             ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1130:41: note: in instantiation of function template specialization 'tt::stl::hash::detail::hash_object<tt::tt_metal::operation::DeviceOperation<>>' requested here
 1130 |     ([&seed](const auto& arg) { seed ^= hash_object(arg) + 0x9e3779b9 + (seed << 6) + (seed >> 2); }(args), ...);
      |                                         ^
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1143:20: note: in instantiation of function template specialization 'tt::stl::hash::detail::hash_objects<int, tt::tt_metal::operation::DeviceOperation<>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>::tensor_args_t>' requested here
 1143 |     return detail::hash_objects(DEFAULT_SEED, args...);
      |                    ^
/home/mmemarian/infra/tt-metal/tt_metal/tools/profiler/op_profiler.hpp:102:31: note: in instantiation of function template specialization 'tt::stl::hash::hash_objects_with_default_seed<int, tt::tt_metal::operation::DeviceOperation<>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>::tensor_args_t>' requested here
  102 |         return tt::stl::hash::hash_objects_with_default_seed(
      |                               ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:286:9: note: in instantiation of function template specialization 'tt::tt_metal::op_profiler::op_meta_data_serialized_json<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>, int, tt::tt_metal::Program, tt::tt_metal::operation::DeviceOperation<>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>::tensor_args_t, std::vector<tt::tt_metal::Tensor>>' requested here
  286 |         TracyOpTTNNDevice(
      |         ^
/home/mmemarian/infra/tt-metal/tt_metal/tools/profiler/op_profiler.hpp:375:43: note: expanded from macro 'TracyOpTTNNDevice'
  375 |     std::string op_message = op_profiler::op_meta_data_serialized_json(                                                \
      |                                           ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:337:5: note: in instantiation of function template specialization 'ttnn::device_operation::detail::launch_on_worker_thread<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>, unsigned char, long, tt::tt_metal::operation::DeviceOperation<>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>::tensor_args_t, std::vector<tt::tt_metal::Tensor>, tt::tt_metal::Device *>' requested here
  337 |     launch_on_worker_thread<device_operation_t>(cq_id, device_operation_id, operation_attributes, tensor_args, tensor_return_value, device);
      |     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:473:32: note: (skipping 10 contexts in backtrace; use -ftemplate-backtrace-limit=0 to see all)
  473 |                 return detail::launch_on_single_device<device_operation_t>(cq_id, operation_attributes, tensor_args);
      |                                ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:642:20: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__base::__visit_alt<std::__variant_detail::__visitation::__variant::__value_visitor<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9)>, const std::__variant_detail::__impl<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
  642 |     return __base::__visit_alt(
      |                    ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:661:12: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__variant::__visit_alt<std::__variant_detail::__visitation::__variant::__value_visitor<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9)>, const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
  661 |     return __visit_alt(
      |            ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:1759:21: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__variant::__visit_value<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9), const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
 1759 |   return __variant::__visit_value(_VSTD::forward<_Visitor>(__visitor),
      |                     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:469:37: note: in instantiation of function template specialization 'std::visit<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9), const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &, void>' requested here
  469 |     auto tensor_return_value = std::visit(
      |                                     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/decorators.hpp:229:48: note: in instantiation of function template specialization 'ttnn::device_operation::detail::invoke<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<tt::tt_metal::Tensor>>>' requested here
  229 |         return ttnn::device_operation::detail::invoke<operation_t>(queue_id, operation_attributes, tensors_args);
      |                                                ^
In file included from <built-in>:1:
In file included from /home/mmemarian/infra/tt-metal/build_Release_tracy/ttnn/CMakeFiles/ttnn.dir/cmake_pch.hxx:5:
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1124:13: error: static assertion failed due to requirement 'tt::stl::concepts::always_false_v<tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>, std::allocator<std::optional<tt::tt_metal::Tensor>>>>>': Type doesn't support hashing using tt::stl::hash::hash_object
 1124 |             tt::stl::concepts::always_false_v<T>, "Type doesn't support hashing using tt::stl::hash::hash_object");
      |             ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1130:41: note: in instantiation of function template specialization 'tt::stl::hash::detail::hash_object<tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>>' requested here
 1130 |     ([&seed](const auto& arg) { seed ^= hash_object(arg) + 0x9e3779b9 + (seed << 6) + (seed >> 2); }(args), ...);
      |                                         ^
/home/mmemarian/infra/tt-metal/tt_metal/tt_stl/reflection.hpp:1143:20: note: in instantiation of function template specialization 'tt::stl::hash::detail::hash_objects<int, tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>::tensor_args_t>' requested here
 1143 |     return detail::hash_objects(DEFAULT_SEED, args...);
      |                    ^
/home/mmemarian/infra/tt-metal/tt_metal/tools/profiler/op_profiler.hpp:102:31: note: in instantiation of function template specialization 'tt::stl::hash::hash_objects_with_default_seed<int, tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>::tensor_args_t>' requested here
  102 |         return tt::stl::hash::hash_objects_with_default_seed(
      |                               ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:286:9: note: in instantiation of function template specialization 'tt::tt_metal::op_profiler::op_meta_data_serialized_json<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, int, tt::tt_metal::Program, tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>::tensor_args_t, std::vector<std::optional<tt::tt_metal::Tensor>>>' requested here
  286 |         TracyOpTTNNDevice(
      |         ^
/home/mmemarian/infra/tt-metal/tt_metal/tools/profiler/op_profiler.hpp:375:43: note: expanded from macro 'TracyOpTTNNDevice'
  375 |     std::string op_message = op_profiler::op_meta_data_serialized_json(                                                \
      |                                           ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:337:5: note: in instantiation of function template specialization 'ttnn::device_operation::detail::launch_on_worker_thread<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, unsigned char, long, tt::tt_metal::operation::DeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>, tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>::tensor_args_t, std::vector<std::optional<tt::tt_metal::Tensor>>, tt::tt_metal::Device *>' requested here
  337 |     launch_on_worker_thread<device_operation_t>(cq_id, device_operation_id, operation_attributes, tensor_args, tensor_return_value, device);
      |     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:473:32: note: (skipping 10 contexts in backtrace; use -ftemplate-backtrace-limit=0 to see all)
  473 |                 return detail::launch_on_single_device<device_operation_t>(cq_id, operation_attributes, tensor_args);
      |                                ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:642:20: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__base::__visit_alt<std::__variant_detail::__visitation::__variant::__value_visitor<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9)>, const std::__variant_detail::__impl<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
  642 |     return __base::__visit_alt(
      |                    ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:661:12: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__variant::__visit_alt<std::__variant_detail::__visitation::__variant::__value_visitor<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9)>, const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
  661 |     return __visit_alt(
      |            ^
/usr/lib/llvm-17/bin/../include/c++/v1/variant:1759:21: note: in instantiation of function template specialization 'std::__variant_detail::__visitation::__variant::__visit_value<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9), const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &>' requested here
 1759 |   return __variant::__visit_value(_VSTD::forward<_Visitor>(__visitor),
      |                     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:469:37: note: in instantiation of function template specialization 'std::visit<(lambda at /home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/device_operation.hpp:470:9), const std::variant<tt::tt_metal::OwnedStorage, tt::tt_metal::DeviceStorage, tt::tt_metal::BorrowedStorage, tt::tt_metal::MultiDeviceHostStorage, tt::tt_metal::MultiDeviceStorage> &, void>' requested here
  469 |     auto tensor_return_value = std::visit(
      |                                     ^
/home/mmemarian/infra/tt-metal/ttnn/cpp/ttnn/decorators.hpp:229:48: note: in instantiation of function template specialization 'ttnn::device_operation::detail::invoke<tt::tt_metal::operation::OldInfraDeviceOperation<std::vector<std::optional<tt::tt_metal::Tensor>>>>' requested here
  229 |         return ttnn::device_operation::detail::invoke<operation_t>(queue_id, operation_attributes, tensors_args);
      |                                                ^
2 errors generated.
[73/75] Building CXX object ttnn/CMakeFiles/ttnn.dir/Unity/unity_56_cxx.cxx.o
ninja: build stopped: subcommand failed.