diff --git a/CMakeLists.txt b/CMakeLists.txt index 0a36f8d106d..a65f2f9b026 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -241,7 +241,6 @@ include(tracy) # Build subdirectories ############################################################################################################################ -add_subdirectory(tt_fabric) add_subdirectory(tt_metal) add_subdirectory(ttnn) @@ -257,15 +256,6 @@ endif() # Install for build artifacts that will upload build/lib include(GNUInstallDirs) -install( - TARGETS - tt_fabric - ARCHIVE - DESTINATION ${CMAKE_INSTALL_LIBDIR} - LIBRARY - DESTINATION ${CMAKE_INSTALL_LIBDIR} - COMPONENT dev -) install( TARGETS tt_metal diff --git a/CODEOWNERS b/CODEOWNERS index e3421407f12..82cb0e61732 100644 --- a/CODEOWNERS +++ b/CODEOWNERS @@ -47,9 +47,6 @@ tests/scripts/t3000/ @tenstorrent/metalium-developers-infra tests/scripts/tg/ @tenstorrent/metalium-developers-infra tests/scripts/tgg/ @tenstorrent/metalium-developers-infra -# fabric -tt_fabric/ @ubcheema @aliuTT @aagarwalTT - # Metalium - public API tt_metal/api @abhullar-tt @pgkeller @aliuTT @tt-aho @tt-dma @tt-asaigal @ubcheema @cfjchu @omilyutin-tt @@ -59,6 +56,9 @@ tt_metal/host_api.hpp @abhullar-tt @pgkeller @aliuTT @tt-aho @tt-dma @tt-asaigal tt_metal/impl/device/ @abhullar-tt @pgkeller @aliuTT @tt-aho @tt-dma @tt-asaigal @ubcheema @davorchap @cfjchu @omilyutin-tt tt_metal/**/requirements*.txt @tenstorrent/metalium-developers-infra +# fabric +tt_metal/fabric/ @ubcheema @aliuTT @aagarwalTT + # metal - dispatch tt_metal/impl/dispatch/kernels/packet_* @ubcheema @aliuTT tt_metal/impl/dispatch/kernels/eth_* @ubcheema @aliuTT diff --git a/tests/tt_metal/tt_fabric/CMakeLists.txt b/tests/tt_metal/tt_fabric/CMakeLists.txt index f18be1886d4..ae7ad23046d 100644 --- a/tests/tt_metal/tt_fabric/CMakeLists.txt +++ b/tests/tt_metal/tt_fabric/CMakeLists.txt @@ -5,7 +5,7 @@ target_link_libraries( fabric_unit_tests PRIVATE tt_metal - tt_fabric + fabric test_common_libs ) @@ -13,7 +13,6 @@ target_include_directories( fabric_unit_tests PRIVATE ${UMD_HOME} - ${PROJECT_SOURCE_DIR}/tt_fabric ${PROJECT_SOURCE_DIR}/tests ${PROJECT_SOURCE_DIR}/tt_metal ${CMAKE_CURRENT_SOURCE_DIR}/common @@ -23,7 +22,7 @@ set_target_properties( fabric_unit_tests PROPERTIES RUNTIME_OUTPUT_DIRECTORY - ${PROJECT_BINARY_DIR}/test/tt_metal/tt_fabric + ${PROJECT_BINARY_DIR}/test/tt_metal/fabric ) gtest_discover_tests(fabric_unit_tests) diff --git a/tests/tt_metal/tt_fabric/fabric_router/test_routing_tables.cpp b/tests/tt_metal/tt_fabric/fabric_router/test_routing_tables.cpp index 9d335001d56..e4ca116fa37 100644 --- a/tests/tt_metal/tt_fabric/fabric_router/test_routing_tables.cpp +++ b/tests/tt_metal/tt_fabric/fabric_router/test_routing_tables.cpp @@ -4,9 +4,9 @@ #include #include "fabric_fixture.hpp" -#include "tt_fabric/control_plane.hpp" -#include "tt_fabric/mesh_graph.hpp" -#include "tt_fabric/routing_table_generator.hpp" +#include "tt_metal/fabric/control_plane.hpp" +#include "tt_metal/fabric/mesh_graph.hpp" +#include "tt_metal/fabric/routing_table_generator.hpp" namespace tt::tt_fabric { namespace fabric_router_tests { @@ -14,21 +14,21 @@ namespace fabric_router_tests { TEST_F(ControlPlaneFixture, TestTGMeshGraphInit) { const std::filesystem::path tg_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; auto mesh_graph_desc = std::make_unique(tg_mesh_graph_desc_path.string()); } TEST_F(ControlPlaneFixture, TestTGControlPlaneInit) { const std::filesystem::path tg_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; auto control_plane = std::make_unique(tg_mesh_graph_desc_path.string()); } TEST_F(ControlPlaneFixture, TestTGFabricRoutes) { const std::filesystem::path tg_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; auto control_plane = std::make_unique(tg_mesh_graph_desc_path.string()); auto valid_chans = control_plane->get_valid_eth_chans_on_routing_plane(0, 0, 3); for (auto chan : valid_chans) { @@ -39,21 +39,21 @@ TEST_F(ControlPlaneFixture, TestTGFabricRoutes) { TEST_F(ControlPlaneFixture, TestT3kMeshGraphInit) { const std::filesystem::path t3k_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; auto mesh_graph_desc = std::make_unique(t3k_mesh_graph_desc_path.string()); } TEST_F(ControlPlaneFixture, TestT3kControlPlaneInit) { const std::filesystem::path t3k_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; auto control_plane = std::make_unique(t3k_mesh_graph_desc_path.string()); } TEST_F(ControlPlaneFixture, TestT3kFabricRoutes) { const std::filesystem::path t3k_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml"; auto control_plane = std::make_unique(t3k_mesh_graph_desc_path.string()); auto valid_chans = control_plane->get_valid_eth_chans_on_routing_plane(0, 0, 0); for (auto chan : valid_chans) { @@ -65,5 +65,12 @@ TEST_F(ControlPlaneFixture, TestT3kFabricRoutes) { } } +TEST_F(ControlPlaneFixture, TestQuantaGalaxyControlPlaneInit) { + const std::filesystem::path quanta_galaxy_mesh_graph_desc_path = + std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / + "tt_metal/fabric/mesh_graph_descriptors/quanta_galaxy_mesh_graph_descriptor.yaml"; + auto control_plane = std::make_unique(quanta_galaxy_mesh_graph_desc_path.string()); +} + } // namespace fabric_router_tests } // namespace tt::tt_fabric diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/CMakeLists.txt b/tests/tt_metal/tt_metal/perf_microbenchmark/CMakeLists.txt index 31e3648d336..df3fb8649a5 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/CMakeLists.txt +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/CMakeLists.txt @@ -65,7 +65,7 @@ foreach(arch ${ARCHITECTURES}) test_metal_common_libs PRIVATE yaml-cpp::yaml-cpp - tt_fabric + fabric ) if(${TEST_SRC} STREQUAL "dispatch/test_pgm_dispatch.cpp") target_link_libraries(${TEST_TARGET} PRIVATE benchmark::benchmark) @@ -76,7 +76,6 @@ foreach(arch ${ARCHITECTURES}) PRIVATE ${PROJECT_SOURCE_DIR}/tt_metal/hw/inc/${arch} "$" - ${PROJECT_SOURCE_DIR}/tt_fabric ${PROJECT_SOURCE_DIR}/ttnn/cpp/ttnn/deprecated # this all should go away and be replaced with link to ttnn ${PROJECT_SOURCE_DIR}/tests ${CMAKE_CURRENT_SOURCE_DIR} diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_rx.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_rx.cpp index efdb7aa794c..346101b8f31 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_rx.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_rx.cpp @@ -5,10 +5,10 @@ // clang-format off #include "debug/dprint.h" #include "dataflow_api.h" -#include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_metal/fabric/hw/inc/tt_fabric.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/common/kernel_utils.hpp" // clang-format on diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_tx.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_tx.cpp index 6abfd6fd57e..731fa976385 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_tx.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_tx.cpp @@ -5,10 +5,10 @@ // clang-format off #include "dataflow_api.h" #include "debug/dprint.h" -#include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_metal/fabric/hw/inc/tt_fabric.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" -#include "tt_fabric/hw/inc/tt_fabric_api.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_api.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/common/kernel_utils.hpp" // clang-format on diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_multi_hop_sanity.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_multi_hop_sanity.cpp index 096370e0c1b..5b53e7b5f87 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_multi_hop_sanity.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_multi_hop_sanity.cpp @@ -6,13 +6,13 @@ #include #include #include -#include "tt_fabric/control_plane.hpp" +#include "tt_metal/fabric/control_plane.hpp" // #include // #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" #include "kernels/tt_fabric_traffic_gen_test.hpp" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" #include "eth_l1_address_map.h" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" using std::vector; using namespace tt; @@ -235,7 +235,7 @@ int main(int argc, char** argv) { try { const std::filesystem::path tg_mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; + "tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml"; auto control_plane = std::make_unique(tg_mesh_graph_desc_path.string()); int num_devices = tt_metal::GetNumAvailableDevices(); @@ -362,7 +362,7 @@ int main(int argc, char** argv) { for (auto logical_core : device_router_cores) { auto router_kernel = tt_metal::CreateKernel( program_map[device.first], - "tt_fabric/impl/kernels/tt_fabric_router.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_router.cpp", logical_core, tt_metal::EthernetConfig{ .noc = tt_metal::NOC::NOC_0, .compile_args = router_compile_args, .defines = defines}); @@ -393,7 +393,7 @@ int main(int argc, char** argv) { auto kernel = tt_metal::CreateKernel( program_map[device.first], - "tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp", {gk_core}, tt_metal::DataMovementConfig{ .processor = tt_metal::DataMovementProcessor::RISCV_0, diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_sanity.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_sanity.cpp index d1d56a9fa53..b9fb71ad80b 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_sanity.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_sanity.cpp @@ -6,14 +6,14 @@ #include #include #include -#include "tt_fabric/control_plane.hpp" -#include "tt_fabric/mesh_graph.hpp" +#include "tt_metal/fabric/mesh_graph.hpp" +#include "tt_metal/fabric/control_plane.hpp" //#include //#include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" #include "kernels/tt_fabric_traffic_gen_test.hpp" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" #include "eth_l1_address_map.h" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" #include #include #include @@ -140,7 +140,7 @@ typedef struct test_board { try { const std::filesystem::path mesh_graph_desc_path = std::filesystem::path(tt::llrt::RunTimeOptions::get_instance().get_root_dir()) / - "tt_fabric/mesh_graph_descriptors" / mesh_graph_descriptor; + "tt_metal/fabric/mesh_graph_descriptors" / mesh_graph_descriptor; control_plane = std::make_unique(mesh_graph_desc_path.string()); } catch (const std::exception& e) { log_fatal(e.what()); @@ -425,7 +425,7 @@ typedef struct test_device { auto kernel = tt_metal::CreateKernel( program_handle, - "tt_fabric/impl/kernels/tt_fabric_router.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_router.cpp", router_logical_cores[i], tt_metal::EthernetConfig{ .noc = tt_metal::NOC::NOC_0, .compile_args = compile_args, .defines = defines}); @@ -451,7 +451,7 @@ typedef struct test_device { if (run_gk_on_idle_ethernet) { kernel = tt_metal::CreateKernel( program_handle, - "tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp", {gk_logical_core}, tt_metal::EthernetConfig{ .eth_mode = Eth::IDLE, @@ -461,7 +461,7 @@ typedef struct test_device { } else { kernel = tt_metal::CreateKernel( program_handle, - "tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp", {gk_logical_core}, tt_metal::DataMovementConfig{ .processor = tt_metal::DataMovementProcessor::RISCV_0, diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_socket_sanity.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_socket_sanity.cpp index e166f43706d..90e6f5cd76c 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_socket_sanity.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tt_fabric_socket_sanity.cpp @@ -6,13 +6,13 @@ #include #include #include -#include "tt_fabric/control_plane.hpp" +#include "tt_metal/fabric/control_plane.hpp" // #include "tt_metal/impl/dispatch/cq_commands.hpp" // #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" #include "kernels/tt_fabric_traffic_gen_test.hpp" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" #include "eth_l1_address_map.h" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" using std::vector; using namespace tt; @@ -391,7 +391,7 @@ int main(int argc, char** argv) { auto kernel = tt_metal::CreateKernel( program_map[device.first], - "tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp", + "tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp", {gk_core}, tt_metal::DataMovementConfig{ .processor = tt_metal::DataMovementProcessor::RISCV_0, diff --git a/tt_fabric/control_plane.hpp b/tt_fabric/control_plane.hpp deleted file mode 100644 index e9faa1377c3..00000000000 --- a/tt_fabric/control_plane.hpp +++ /dev/null @@ -1,75 +0,0 @@ -// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc. -// -// SPDX-License-Identifier: Apache-2.0 - -#pragma once - -#include "routing_table_generator.hpp" -#include -#include -#include "hw/inc/routing_table.h" - -namespace tt::tt_fabric { - -class ControlPlane { - public: - explicit ControlPlane(const std::string& mesh_graph_desc_yaml_file); - ~ControlPlane() = default; - void initialize_from_mesh_graph_desc_file(const std::string& mesh_graph_desc_file); - - // Takes RoutingTableGenerator table and converts to routing tables for each ethernet port - void convert_fabric_routing_table_to_chip_routing_table(); - - void write_routing_tables_to_chip(mesh_id_t mesh_id, chip_id_t chip_id) const; - void configure_routing_tables() const; - - // Printing functions - void print_routing_tables() const; - void print_ethernet_channels() const; - - // Return mesh_id, chip_id from physical chip id - std::pair get_mesh_chip_id_from_physical_chip_id(chip_id_t physical_chip_id) const; - chip_id_t get_physical_chip_id_from_mesh_chip_id(const std::pair& mesh_chip_id) const; - - // Return valid ethernet channels on the specificed routing plane - std::vector get_valid_eth_chans_on_routing_plane( - mesh_id_t mesh_id, chip_id_t chip_id, routing_plane_id_t routing_plane_id) const; - - // Return path from device to device in the fabric - std::vector> get_fabric_route( - mesh_id_t src_mesh_id, - chip_id_t src_chip_id, - mesh_id_t dst_mesh_id, - chip_id_t dst_chip_id, - chan_id_t src_chan_id) const; - - private: - std::unique_ptr routing_table_generator_; - std::vector> logical_mesh_chip_id_to_physical_chip_id_mapping_; - // map[mesh_id][chip_id][direction] has a list of ethernet channels in that direction - std::vector>>> - router_port_directions_to_physical_eth_chan_map_; - // tables[mesh_id][chip_id][eth_chan] - std::vector>>> - intra_mesh_routing_tables_; // table that will be written to each ethernet core - std::vector>>> - inter_mesh_routing_tables_; // table that will be written to each ethernet core - - // Tries to get a valid downstream channel from the candidate_target_chans - // First along same routing plane, but if not available, take round robin from candidates - chan_id_t get_downstream_eth_chan_id( - chan_id_t src_chan_id, const std::vector& candidate_target_chans) const; - - std::vector get_mesh_physical_chip_ids( - std::uint32_t mesh_ns_size, - std::uint32_t mesh_ew_size, - std::uint32_t num_ports_per_side, - std::uint32_t nw_chip_physical_chip_id); - - std::tuple get_connected_mesh_chip_chan_ids( - mesh_id_t mesh_id, chip_id_t chip_id, chan_id_t chan_id) const; - - routing_plane_id_t get_routing_plane_id(chan_id_t eth_chan_id) const; -}; - -} // namespace tt::tt_fabric diff --git a/tt_fabric/routing_table_generator.hpp b/tt_fabric/routing_table_generator.hpp deleted file mode 100644 index 0034ad05a0d..00000000000 --- a/tt_fabric/routing_table_generator.hpp +++ /dev/null @@ -1,60 +0,0 @@ -// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc. -// -// SPDX-License-Identifier: Apache-2.0 - -#pragma once -#include -#include "mesh_graph.hpp" - -namespace tt::tt_fabric { - -using RoutingTable = - std::vector>>; // [mesh_id][chip_id][target_chip_or_mesh_id] - -class RoutingTableGenerator { - public: - explicit RoutingTableGenerator(const std::string& mesh_graph_desc_yaml_file); - ~RoutingTableGenerator() = default; - - void dump_to_yaml(); - void load_from_yaml(); - - void print_connectivity() const { this->mesh_graph_->print_connectivity(); } - - const IntraMeshConnectivity& get_intra_mesh_connectivity() const { - return this->mesh_graph_->get_intra_mesh_connectivity(); - } - const InterMeshConnectivity& get_inter_mesh_connectivity() const { - return this->mesh_graph_->get_inter_mesh_connectivity(); - } - const ChipSpec& get_chip_spec() const { return this->mesh_graph_->get_chip_spec(); } - - std::uint32_t get_mesh_ns_size(mesh_id_t mesh_id) const { return this->mesh_graph_->get_mesh_ns_size(mesh_id); } - std::uint32_t get_mesh_ew_size(mesh_id_t mesh_id) const { return this->mesh_graph_->get_mesh_ew_size(mesh_id); } - - RoutingTable get_intra_mesh_table() const { return this->intra_mesh_table_; } - RoutingTable get_inter_mesh_table() const { return this->inter_mesh_table_; } - - void print_routing_tables() const; - - private: - std::unique_ptr mesh_graph_; - // configurable in future architectures - const uint32_t max_nodes_in_mesh_ = 1024; - const uint32_t max_num_meshes_ = 1024; - - std::vector mesh_sizes; - - RoutingTable intra_mesh_table_; - RoutingTable inter_mesh_table_; - - std::vector>>> get_paths_to_all_meshes( - mesh_id_t src, const InterMeshConnectivity& inter_mesh_connectivity); - void generate_intramesh_routing_table(const IntraMeshConnectivity& intra_mesh_connectivity); - // when generating intermesh routing table, we use the intramesh connectivity table to find the shortest path to - // the exit chip - void generate_intermesh_routing_table( - const InterMeshConnectivity& inter_mesh_connectivity, const IntraMeshConnectivity& intra_mesh_connectivity); -}; - -} // namespace tt::tt_fabric diff --git a/tt_metal/CMakeLists.txt b/tt_metal/CMakeLists.txt index 768c9318eac..2d314c105d0 100644 --- a/tt_metal/CMakeLists.txt +++ b/tt_metal/CMakeLists.txt @@ -30,6 +30,7 @@ target_link_libraries( llrt detail distributed + fabric HAL::grayskull HAL::wormhole HAL::blackhole @@ -87,3 +88,4 @@ add_subdirectory(impl) add_subdirectory(detail) add_subdirectory(distributed) add_subdirectory(tt_stl) +add_subdirectory(fabric) diff --git a/tt_fabric/CMakeLists.txt b/tt_metal/fabric/CMakeLists.txt similarity index 64% rename from tt_fabric/CMakeLists.txt rename to tt_metal/fabric/CMakeLists.txt index 34add9c0350..f6cf8109a3e 100644 --- a/tt_fabric/CMakeLists.txt +++ b/tt_metal/fabric/CMakeLists.txt @@ -1,29 +1,29 @@ -add_library(tt_fabric) -add_library(TT::Fabric ALIAS tt_fabric) +add_library(fabric) +add_library(TT::Fabric ALIAS fabric) target_sources( - tt_fabric + fabric PRIVATE control_plane.cpp routing_table_generator.cpp mesh_graph.cpp ) -target_include_directories(tt_fabric PRIVATE .) +target_include_directories(fabric PRIVATE .) target_link_libraries( - tt_fabric + fabric PRIVATE - Metalium::Metal umd::device - metal_common_libs magic_enum fmt::fmt-header-only yaml-cpp::yaml-cpp + Metalium::Metal::Impl + TT::Metalium::HostDevCommon ) target_precompile_headers( - tt_fabric + fabric PRIVATE @@ -33,10 +33,10 @@ target_precompile_headers( ) -target_compile_options(tt_fabric PRIVATE -Wno-int-to-pointer-cast) +target_compile_options(fabric PRIVATE -Wno-int-to-pointer-cast) set_target_properties( - tt_fabric + fabric PROPERTIES INSTALL_RPATH "${PROJECT_BINARY_DIR}/lib" diff --git a/tt_fabric/control_plane.cpp b/tt_metal/fabric/control_plane.cpp similarity index 100% rename from tt_fabric/control_plane.cpp rename to tt_metal/fabric/control_plane.cpp diff --git a/tt_metal/fabric/control_plane.hpp b/tt_metal/fabric/control_plane.hpp new file mode 100644 index 00000000000..14c62cdcbf1 --- /dev/null +++ b/tt_metal/fabric/control_plane.hpp @@ -0,0 +1,76 @@ +// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +#include "routing_table_generator.hpp" +#include "fabric/hw/inc/routing_table.h" + +namespace tt::tt_fabric { + +class ControlPlane { +public: + explicit ControlPlane(const std::string& mesh_graph_desc_yaml_file); + ~ControlPlane() = default; + void initialize_from_mesh_graph_desc_file(const std::string& mesh_graph_desc_file); + + // Takes RoutingTableGenerator table and converts to routing tables for each ethernet port + void convert_fabric_routing_table_to_chip_routing_table(); + + void write_routing_tables_to_chip(mesh_id_t mesh_id, chip_id_t chip_id) const; + void configure_routing_tables() const; + + // Printing functions + void print_routing_tables() const; + void print_ethernet_channels() const; + + // Return mesh_id, chip_id from physical chip id + std::pair get_mesh_chip_id_from_physical_chip_id(chip_id_t physical_chip_id) const; + chip_id_t get_physical_chip_id_from_mesh_chip_id(const std::pair& mesh_chip_id) const; + + // Return valid ethernet channels on the specificed routing plane + std::vector get_valid_eth_chans_on_routing_plane( + mesh_id_t mesh_id, chip_id_t chip_id, routing_plane_id_t routing_plane_id) const; + + // Return path from device to device in the fabric + std::vector> get_fabric_route( + mesh_id_t src_mesh_id, + chip_id_t src_chip_id, + mesh_id_t dst_mesh_id, + chip_id_t dst_chip_id, + chan_id_t src_chan_id) const; + +private: + std::unique_ptr routing_table_generator_; + std::vector> logical_mesh_chip_id_to_physical_chip_id_mapping_; + // map[mesh_id][chip_id][direction] has a list of ethernet channels in that direction + std::vector>>> + router_port_directions_to_physical_eth_chan_map_; + // tables[mesh_id][chip_id][eth_chan] + std::vector>>> + intra_mesh_routing_tables_; // table that will be written to each ethernet core + std::vector>>> + inter_mesh_routing_tables_; // table that will be written to each ethernet core + + // Tries to get a valid downstream channel from the candidate_target_chans + // First along same routing plane, but if not available, take round robin from candidates + chan_id_t get_downstream_eth_chan_id( + chan_id_t src_chan_id, const std::vector& candidate_target_chans) const; + + std::vector get_mesh_physical_chip_ids( + std::uint32_t mesh_ns_size, + std::uint32_t mesh_ew_size, + std::uint32_t num_ports_per_side, + std::uint32_t nw_chip_physical_chip_id); + + std::tuple get_connected_mesh_chip_chan_ids( + mesh_id_t mesh_id, chip_id_t chip_id, chan_id_t chan_id) const; + + routing_plane_id_t get_routing_plane_id(chan_id_t eth_chan_id) const; +}; + +} // namespace tt::tt_fabric diff --git a/tt_fabric/hw/inc/eth_chan_noc_mapping.h b/tt_metal/fabric/hw/inc/eth_chan_noc_mapping.h similarity index 100% rename from tt_fabric/hw/inc/eth_chan_noc_mapping.h rename to tt_metal/fabric/hw/inc/eth_chan_noc_mapping.h diff --git a/tt_fabric/hw/inc/routing_table.h b/tt_metal/fabric/hw/inc/routing_table.h similarity index 88% rename from tt_fabric/hw/inc/routing_table.h rename to tt_metal/fabric/hw/inc/routing_table.h index 70c862cc009..2c24c76401c 100644 --- a/tt_fabric/hw/inc/routing_table.h +++ b/tt_metal/fabric/hw/inc/routing_table.h @@ -32,7 +32,9 @@ static constexpr std::uint32_t MODULO_LOG_BASE_2 = (1 << LOG_BASE_2_NUM_CHANNELS static constexpr std::uint32_t NUM_TABLE_ENTRIES = MAX_MESH_SIZE >> LOG_BASE_2_NUM_CHANNELS_PER_UINT32; static_assert(MAX_MESH_SIZE == MAX_NUM_MESHES, "MAX_MESH_SIZE must be equal to MAX_NUM_MESHES"); -static_assert((sizeof(std::uint32_t) / sizeof(chan_id_t)) == NUM_CHANNELS_PER_UINT32, "LOG_BASE_2_NUM_CHANNELS_PER_UINT32 must be equal to log2(sizeof(std::uint32_t) / sizeof(chan_id_t))"); +static_assert( + (sizeof(std::uint32_t) / sizeof(chan_id_t)) == NUM_CHANNELS_PER_UINT32, + "LOG_BASE_2_NUM_CHANNELS_PER_UINT32 must be equal to log2(sizeof(std::uint32_t) / sizeof(chan_id_t))"); enum eth_chan_magic_values { INVALID_DIRECTION = 0xDD, @@ -40,14 +42,14 @@ enum eth_chan_magic_values { }; struct routing_table_t { - chan_id_t dest_entry[MAX_MESH_SIZE]; + chan_id_t dest_entry[MAX_MESH_SIZE]; }; struct port_direction_t { - chan_id_t north; - chan_id_t south; - chan_id_t east; - chan_id_t west; + chan_id_t north; + chan_id_t south; + chan_id_t east; + chan_id_t west; }; struct fabric_router_l1_config_t { diff --git a/tt_fabric/hw/inc/tt_fabric.h b/tt_metal/fabric/hw/inc/tt_fabric.h similarity index 99% rename from tt_fabric/hw/inc/tt_fabric.h rename to tt_metal/fabric/hw/inc/tt_fabric.h index 52944d8c8c2..084193eab90 100644 --- a/tt_fabric/hw/inc/tt_fabric.h +++ b/tt_metal/fabric/hw/inc/tt_fabric.h @@ -9,9 +9,9 @@ #include "dataflow_api.h" #include "noc_overlay_parameters.h" #include "ethernet/dataflow_api.h" -#include "tt_fabric/hw/inc/routing_table.h" -#include "tt_fabric/hw/inc/tt_fabric_interface.h" -#include "tt_fabric/hw/inc/eth_chan_noc_mapping.h" +#include "tt_metal/fabric/hw/inc/routing_table.h" +#include "tt_metal/fabric/hw/inc/tt_fabric_interface.h" +#include "tt_metal/fabric/hw/inc/eth_chan_noc_mapping.h" using namespace tt::tt_fabric; @@ -423,7 +423,7 @@ typedef struct fvc_producer_state { inline void advance_next_packet() { if (this->get_num_words_available() >= PACKET_HEADER_SIZE_WORDS) { tt_l1_ptr uint32_t* packet_header_ptr = (uint32_t*)¤t_packet_header; - tt_l1_ptr volatile uint32_t* next_header_ptr = + volatile tt_l1_ptr uint32_t* next_header_ptr = reinterpret_cast(get_local_buffer_read_addr()); uint32_t words_before_wrap = words_before_buffer_wrap(fvc_out_rdptr); uint32_t dwords_to_copy = PACKET_HEADER_SIZE_BYTES / 4; diff --git a/tt_fabric/hw/inc/tt_fabric_api.h b/tt_metal/fabric/hw/inc/tt_fabric_api.h similarity index 100% rename from tt_fabric/hw/inc/tt_fabric_api.h rename to tt_metal/fabric/hw/inc/tt_fabric_api.h diff --git a/tt_fabric/hw/inc/tt_fabric_interface.h b/tt_metal/fabric/hw/inc/tt_fabric_interface.h similarity index 100% rename from tt_fabric/hw/inc/tt_fabric_interface.h rename to tt_metal/fabric/hw/inc/tt_fabric_interface.h diff --git a/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp b/tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp similarity index 99% rename from tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp rename to tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp index b90892d5e5b..0ccc74f6bc0 100644 --- a/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp +++ b/tt_metal/fabric/impl/kernels/tt_fabric_gatekeeper.cpp @@ -4,7 +4,7 @@ // clang-format off #include "dataflow_api.h" -#include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_metal/fabric/hw/inc/tt_fabric.h" #include "debug/dprint.h" // clang-format on diff --git a/tt_fabric/impl/kernels/tt_fabric_router.cpp b/tt_metal/fabric/impl/kernels/tt_fabric_router.cpp similarity index 99% rename from tt_fabric/impl/kernels/tt_fabric_router.cpp rename to tt_metal/fabric/impl/kernels/tt_fabric_router.cpp index efaa73cc58c..9235b9ba800 100644 --- a/tt_fabric/impl/kernels/tt_fabric_router.cpp +++ b/tt_metal/fabric/impl/kernels/tt_fabric_router.cpp @@ -4,7 +4,7 @@ // clang-format off #include "dataflow_api.h" -#include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_metal/fabric/hw/inc/tt_fabric.h" // clang-format on using namespace tt::tt_fabric; @@ -43,7 +43,7 @@ constexpr uint32_t PQ_TEST_MISC_INDEX = 16; // careful, may be null tt_l1_ptr uint32_t* const kernel_status = reinterpret_cast(kernel_status_buf_addr_arg); -tt_l1_ptr volatile chan_req_buf* fvc_consumer_req_buf = +volatile tt_l1_ptr chan_req_buf* fvc_consumer_req_buf = reinterpret_cast(FABRIC_ROUTER_REQ_QUEUE_START); volatile tt_l1_ptr fabric_router_l1_config_t* routing_table = reinterpret_cast(eth_l1_mem::address_map::FABRIC_ROUTER_CONFIG_BASE); diff --git a/tt_fabric/mesh_graph.cpp b/tt_metal/fabric/mesh_graph.cpp similarity index 100% rename from tt_fabric/mesh_graph.cpp rename to tt_metal/fabric/mesh_graph.cpp diff --git a/tt_fabric/mesh_graph.hpp b/tt_metal/fabric/mesh_graph.hpp similarity index 54% rename from tt_fabric/mesh_graph.hpp rename to tt_metal/fabric/mesh_graph.hpp index 414b8947527..950324b79cf 100644 --- a/tt_fabric/mesh_graph.hpp +++ b/tt_metal/fabric/mesh_graph.hpp @@ -55,39 +55,36 @@ using InterMeshConnectivity = std::vector>>; class MeshGraph { - public: - explicit MeshGraph(const std::string& mesh_graph_desc_file_path); - MeshGraph() = delete; - ~MeshGraph() = default; - - void print_connectivity() const; - - const IntraMeshConnectivity& get_intra_mesh_connectivity() const { return intra_mesh_connectivity_; } - const InterMeshConnectivity& get_inter_mesh_connectivity() const { return inter_mesh_connectivity_; } - - const ChipSpec& get_chip_spec() const { return chip_spec_; } - - std::uint32_t get_mesh_ns_size(mesh_id_t mesh_id) const { return mesh_shapes_[mesh_id].first; } - std::uint32_t get_mesh_ew_size(mesh_id_t mesh_id) const { return mesh_shapes_[mesh_id].second; } - - private: - std::unordered_map get_valid_connections( - chip_id_t src_chip_id, - std::uint32_t row_size, - std::uint32_t num_chips_in_mesh, - FabricType fabric_type) const; - void initialize_from_yaml(const std::string& mesh_graph_desc_file_path); - - void add_to_connectivity( - mesh_id_t src_mesh_id, - chip_id_t src_chip_id, - chip_id_t dest_mesh_id, - chip_id_t dest_chip_id, - RoutingDirection port_direction); - - ChipSpec chip_spec_; - std::vector> mesh_shapes_; - IntraMeshConnectivity intra_mesh_connectivity_; - InterMeshConnectivity inter_mesh_connectivity_; +public: + explicit MeshGraph(const std::string& mesh_graph_desc_file_path); + MeshGraph() = delete; + ~MeshGraph() = default; + + void print_connectivity() const; + + const IntraMeshConnectivity& get_intra_mesh_connectivity() const { return intra_mesh_connectivity_; } + const InterMeshConnectivity& get_inter_mesh_connectivity() const { return inter_mesh_connectivity_; } + + const ChipSpec& get_chip_spec() const { return chip_spec_; } + + std::uint32_t get_mesh_ns_size(mesh_id_t mesh_id) const { return mesh_shapes_[mesh_id].first; } + std::uint32_t get_mesh_ew_size(mesh_id_t mesh_id) const { return mesh_shapes_[mesh_id].second; } + +private: + std::unordered_map get_valid_connections( + chip_id_t src_chip_id, std::uint32_t row_size, std::uint32_t num_chips_in_mesh, FabricType fabric_type) const; + void initialize_from_yaml(const std::string& mesh_graph_desc_file_path); + + void add_to_connectivity( + mesh_id_t src_mesh_id, + chip_id_t src_chip_id, + chip_id_t dest_mesh_id, + chip_id_t dest_chip_id, + RoutingDirection port_direction); + + ChipSpec chip_spec_; + std::vector> mesh_shapes_; + IntraMeshConnectivity intra_mesh_connectivity_; + InterMeshConnectivity inter_mesh_connectivity_; }; } // namespace tt::tt_fabric diff --git a/tt_fabric/mesh_graph_descriptors/n300_mesh_graph_descriptor.yaml b/tt_metal/fabric/mesh_graph_descriptors/n300_mesh_graph_descriptor.yaml similarity index 100% rename from tt_fabric/mesh_graph_descriptors/n300_mesh_graph_descriptor.yaml rename to tt_metal/fabric/mesh_graph_descriptors/n300_mesh_graph_descriptor.yaml diff --git a/tt_metal/fabric/mesh_graph_descriptors/quanta_galaxy_mesh_graph_descriptor.yaml b/tt_metal/fabric/mesh_graph_descriptors/quanta_galaxy_mesh_graph_descriptor.yaml new file mode 100644 index 00000000000..72aa6e210c9 --- /dev/null +++ b/tt_metal/fabric/mesh_graph_descriptors/quanta_galaxy_mesh_graph_descriptor.yaml @@ -0,0 +1,27 @@ +ChipSpec: { + arch: wormhole_b0, + ethernet_ports: { + N: 4, + E: 4, + S: 4, + W: 4, + } +} + + +Board: [ + { name: Galaxy, + type: Mesh, + topology: [4, 8]} + ] + +Mesh: [ +{ + id: 0, + board: Galaxy, + topology: [1, 1], + host_mapping: [[]]} +] + +Graph: [ +] diff --git a/tt_fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml b/tt_metal/fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml similarity index 100% rename from tt_fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml rename to tt_metal/fabric/mesh_graph_descriptors/t3k_mesh_graph_descriptor.yaml diff --git a/tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml b/tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml similarity index 100% rename from tt_fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml rename to tt_metal/fabric/mesh_graph_descriptors/tg_mesh_graph_descriptor.yaml diff --git a/tt_fabric/routing_table_generator.cpp b/tt_metal/fabric/routing_table_generator.cpp similarity index 100% rename from tt_fabric/routing_table_generator.cpp rename to tt_metal/fabric/routing_table_generator.cpp diff --git a/tt_metal/fabric/routing_table_generator.hpp b/tt_metal/fabric/routing_table_generator.hpp new file mode 100644 index 00000000000..ac57204ef1e --- /dev/null +++ b/tt_metal/fabric/routing_table_generator.hpp @@ -0,0 +1,60 @@ +// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#pragma once +#include +#include "mesh_graph.hpp" + +namespace tt::tt_fabric { + +using RoutingTable = + std::vector>>; // [mesh_id][chip_id][target_chip_or_mesh_id] + +class RoutingTableGenerator { +public: + explicit RoutingTableGenerator(const std::string& mesh_graph_desc_yaml_file); + ~RoutingTableGenerator() = default; + + void dump_to_yaml(); + void load_from_yaml(); + + void print_connectivity() const { this->mesh_graph_->print_connectivity(); } + + const IntraMeshConnectivity& get_intra_mesh_connectivity() const { + return this->mesh_graph_->get_intra_mesh_connectivity(); + } + const InterMeshConnectivity& get_inter_mesh_connectivity() const { + return this->mesh_graph_->get_inter_mesh_connectivity(); + } + const ChipSpec& get_chip_spec() const { return this->mesh_graph_->get_chip_spec(); } + + std::uint32_t get_mesh_ns_size(mesh_id_t mesh_id) const { return this->mesh_graph_->get_mesh_ns_size(mesh_id); } + std::uint32_t get_mesh_ew_size(mesh_id_t mesh_id) const { return this->mesh_graph_->get_mesh_ew_size(mesh_id); } + + RoutingTable get_intra_mesh_table() const { return this->intra_mesh_table_; } + RoutingTable get_inter_mesh_table() const { return this->inter_mesh_table_; } + + void print_routing_tables() const; + +private: + std::unique_ptr mesh_graph_; + // configurable in future architectures + const uint32_t max_nodes_in_mesh_ = 1024; + const uint32_t max_num_meshes_ = 1024; + + std::vector mesh_sizes; + + RoutingTable intra_mesh_table_; + RoutingTable inter_mesh_table_; + + std::vector>>> get_paths_to_all_meshes( + mesh_id_t src, const InterMeshConnectivity& inter_mesh_connectivity); + void generate_intramesh_routing_table(const IntraMeshConnectivity& intra_mesh_connectivity); + // when generating intermesh routing table, we use the intramesh connectivity table to find the shortest path to + // the exit chip + void generate_intermesh_routing_table( + const InterMeshConnectivity& inter_mesh_connectivity, const IntraMeshConnectivity& intra_mesh_connectivity); +}; + +} // namespace tt::tt_fabric