From 02fb2125f3fe1f05afd38bcea7993ccb7df87313 Mon Sep 17 00:00:00 2001 From: Daiki Aminaka Date: Mon, 10 Feb 2025 21:34:36 -0800 Subject: [PATCH] Refactoring same definitions (#17747) ### Ticket N/A ### Problem description There are same definitions spreading to multiple files. The name is overwrapping with other file's one, so refactoring to make it really unique ### What's changed Fix name - PACKET_QUEUE_TEST to TT_FABRIC_STATUS - PQ_TEST to TT_FABRIC - move common test utilities to test_common.hpp ### Checklist - [ ] [All post commit](https://github.com/tenstorrent/tt-metal/actions/workflows/all-post-commit-workflows.yaml) CI passes - [ ] [Blackhole Post commit](https://github.com/tenstorrent/tt-metal/actions/workflows/blackhole-post-commit.yaml) CI passes (if applicable) - [ ] [Model regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-models.yaml) CI passes (if applicable) - [ ] [Device performance regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-device-models.yaml) CI passes (if applicable) - [ ] **(For models and ops writers)** Full [new models tests](https://github.com/tenstorrent/tt-metal/actions/workflows/full-new-models-suite.yaml) CI passes (if applicable) - [ ] New/Existing tests provide coverage for changes --- .../dispatch/test_prefetcher.cpp | 1 - .../routing/kernels/traffic_gen.hpp | 2 +- .../routing/kernels/traffic_gen_test.hpp | 41 ---------- .../routing/kernels/traffic_gen_tx.cpp | 1 + .../routing/kernels/tt_fabric_traffic_gen.hpp | 2 +- .../kernels/tt_fabric_traffic_gen_rx.cpp | 20 ++--- .../kernels/tt_fabric_traffic_gen_test.hpp | 78 ------------------- .../routing/kernels/tt_fabric_tx_ubench.cpp | 14 ++-- .../routing/test_common.hpp | 9 +++ .../routing/test_mux_demux.cpp | 3 +- .../routing/test_mux_demux_2level.cpp | 2 +- .../test_tt_fabric_multi_hop_sanity.cpp | 26 +++---- .../routing/test_tt_fabric_sanity.cpp | 35 +++++---- .../routing/test_tt_fabric_socket_sanity.cpp | 26 +++---- .../routing/test_tx_rx.cpp | 2 +- .../routing/test_vc_bi_tunnel_2ep.cpp | 3 +- .../routing/test_vc_bi_tunnel_4ep.cpp | 3 +- .../routing/test_vc_loopback_tunnel.cpp | 3 +- .../routing/test_vc_mux_demux.cpp | 3 +- .../routing/test_vc_uni_tunnel.cpp | 3 +- tt_fabric/hw/inc/tt_fabric_status.h | 45 +++++++++++ .../impl/kernels/tt_fabric_gatekeeper.cpp | 35 +++------ tt_fabric/impl/kernels/tt_fabric_router.cpp | 49 +++++------- .../dispatch/kernels/packet_queue_ctrl.hpp | 11 +++ 24 files changed, 162 insertions(+), 255 deletions(-) delete mode 100644 tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp delete mode 100644 tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp create mode 100644 tt_fabric/hw/inc/tt_fabric_status.h diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp index ab2483709e2..0b1dc88bec3 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp @@ -17,7 +17,6 @@ #include "common.h" #include "tt_cluster.hpp" #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp" #include #include "llrt.hpp" diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen.hpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen.hpp index a255f46c798..01b9dedaae2 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen.hpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen.hpp @@ -5,7 +5,7 @@ #pragma once #include "debug/dprint.h" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp" +#include "tt_fabric/hw/inc/tt_fabric_status.h" inline uint32_t prng_next(uint32_t n) { uint32_t x = n; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp deleted file mode 100644 index 6e28268ef98..00000000000 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp +++ /dev/null @@ -1,41 +0,0 @@ -// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc. -// -// SPDX-License-Identifier: Apache-2.0 - -#pragma once - -#include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include -#include - -inline const char* packet_queue_test_status_to_string(uint32_t status) { - switch (status) { - case PACKET_QUEUE_TEST_STARTED: return "STARTED"; - case PACKET_QUEUE_TEST_PASS: return "DONE/OK"; - case PACKET_QUEUE_TEST_TIMEOUT: return "TIMEOUT"; - case PACKET_QUEUE_TEST_DATA_MISMATCH: return "DATA_MISMATCH"; - default: return "UNKNOWN"; - } -} - -inline uint64_t get_64b_result(uint32_t* buf, uint32_t index) { - return (((uint64_t)buf[index]) << 32) | buf[index + 1]; -} - -inline uint64_t get_64b_result(const std::vector& vec, uint32_t index) { - return (((uint64_t)vec[index]) << 32) | vec[index + 1]; -} - -#define TX_TEST_IDX_TOT_DATA_WORDS PQ_TEST_MISC_INDEX + 1 -#define TX_TEST_IDX_NPKT PQ_TEST_MISC_INDEX + 3 -#define TX_TEST_IDX_WORDS_FLUSHED PQ_TEST_MISC_INDEX + 5 -#define TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 7 -#define TX_TEST_IDX_MANY_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 9 -#define TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 11 -// #define TX_TEST_IDX_ PQ_TEST_MISC_INDEX + -// #define TX_TEST_IDX_ PQ_TEST_MISC_INDEX + - -enum class pkt_dest_size_choices_t { - RANDOM = 0, - SAME_START_RNDROBIN_FIX_SIZE = 1 // max packet size used -}; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_tx.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_tx.cpp index 24a7decd1bd..57812ccde36 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_tx.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_tx.cpp @@ -5,6 +5,7 @@ #include "dataflow_api.h" #include "debug/dprint.h" #include "tt_metal/impl/dispatch/kernels/packet_queue.hpp" +#include "tt_fabric/hw/inc/tt_fabric_status.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen.hpp" constexpr uint32_t src_endpoint_id = get_compile_time_arg_val(0); diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp index 23a32149192..19fcdc79dbd 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen.hpp @@ -5,7 +5,7 @@ #pragma once #include "debug/dprint.h" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp" +#include "tt_fabric/hw/inc/tt_fabric_status.h" #define is_power_of_2(x) (((x) > 0) && (((x) & ((x) - 1)) == 0)) 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..4c29d8b4ef9 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 @@ -7,7 +7,7 @@ #include "dataflow_api.h" #include "tt_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_status.h" #include "tt_fabric/hw/inc/tt_fabric_interface.h" #include "tests/tt_metal/tt_metal/perf_microbenchmark/common/kernel_utils.hpp" // clang-format on @@ -61,8 +61,8 @@ void kernel_main() { rx_buf_size = get_arg_val(increment_arg_idx(rt_args_idx)); zero_l1_buf(test_results, test_results_size_bytes); - test_results[PQ_TEST_STATUS_INDEX] = PACKET_QUEUE_TEST_STARTED; - test_results[PQ_TEST_MISC_INDEX] = 0xff000000; + test_results[TT_FABRIC_STATUS_INDEX] = TT_FABRIC_STATUS_STARTED; + test_results[TT_FABRIC_MISC_INDEX] = 0xff000000; if constexpr (ASYNC_WR & test_command) { uint32_t packet_rnd_seed; @@ -174,9 +174,9 @@ void kernel_main() { read_addr, curr_payload_words, start_val, mismatch_addr, mismatch_val, expected_val); if (!match) { async_wr_check_failed = true; - test_results[PQ_TEST_MISC_INDEX + 12] = mismatch_addr; - test_results[PQ_TEST_MISC_INDEX + 13] = mismatch_val; - test_results[PQ_TEST_MISC_INDEX + 14] = expected_val; + test_results[TT_FABRIC_MISC_INDEX + 12] = mismatch_addr; + test_results[TT_FABRIC_MISC_INDEX + 13] = mismatch_val; + test_results[TT_FABRIC_MISC_INDEX + 14] = expected_val; break; } } @@ -200,13 +200,13 @@ void kernel_main() { } // write out results - set_64b_result(test_results, processed_packet_words, PQ_TEST_WORD_CNT_INDEX); + set_64b_result(test_results, processed_packet_words, TT_FABRIC_WORD_CNT_INDEX); set_64b_result(test_results, num_packets, TX_TEST_IDX_NPKT); if (async_wr_check_failed) { - test_results[PQ_TEST_STATUS_INDEX] = PACKET_QUEUE_TEST_DATA_MISMATCH; + test_results[TT_FABRIC_STATUS_INDEX] = TT_FABRIC_STATUS_DATA_MISMATCH; } else { - test_results[PQ_TEST_STATUS_INDEX] = PACKET_QUEUE_TEST_PASS; - test_results[PQ_TEST_MISC_INDEX] = 0xff000005; + test_results[TT_FABRIC_STATUS_INDEX] = TT_FABRIC_STATUS_PASS; + test_results[TT_FABRIC_MISC_INDEX] = 0xff000005; } } diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp deleted file mode 100644 index ac4ebaee8e3..00000000000 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_traffic_gen_test.hpp +++ /dev/null @@ -1,78 +0,0 @@ -// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc. -// -// SPDX-License-Identifier: Apache-2.0 - -#pragma once - -//#include "tt_metal/impl/dispatch/kernels/tt_fabric.hpp" -#include -#include - -constexpr uint32_t PACKET_QUEUE_STAUS_MASK = 0xabc00000; -constexpr uint32_t PACKET_QUEUE_TEST_STARTED = PACKET_QUEUE_STAUS_MASK | 0x0; -constexpr uint32_t PACKET_QUEUE_TEST_PASS = PACKET_QUEUE_STAUS_MASK | 0x1; -constexpr uint32_t PACKET_QUEUE_TEST_TIMEOUT = PACKET_QUEUE_STAUS_MASK | 0xdead0; -constexpr uint32_t PACKET_QUEUE_TEST_BAD_HEADER = PACKET_QUEUE_STAUS_MASK | 0xdead1; -constexpr uint32_t PACKET_QUEUE_TEST_DATA_MISMATCH = PACKET_QUEUE_STAUS_MASK | 0x3; - -// indexes of return values in test results buffer -constexpr uint32_t PQ_TEST_STATUS_INDEX = 0; -constexpr uint32_t PQ_TEST_WORD_CNT_INDEX = 2; -constexpr uint32_t PQ_TEST_CYCLES_INDEX = 4; -constexpr uint32_t PQ_TEST_ITER_INDEX = 6; -constexpr uint32_t PQ_TEST_MISC_INDEX = 16; - -/* -inline const char *packet_queue_test_status_to_string(uint32_t status) { - switch (status) { - case TT_FABRIC_TEST_STARTED: - return "STARTED"; - case TT_FABRIC_TEST_PASS: - return "DONE/OK"; - case TT_FABRIC_TEST_TIMEOUT: - return "TIMEOUT"; - case TT_FABRIC_TEST_DATA_MISMATCH: - return "DATA_MISMATCH"; - default: - return "UNKNOWN"; - } -} -*/ - -inline const char *packet_queue_test_status_to_string(uint32_t status) { - switch (status) { - case PACKET_QUEUE_TEST_STARTED: - return "STARTED"; - case PACKET_QUEUE_TEST_PASS: - return "DONE/OK"; - case PACKET_QUEUE_TEST_TIMEOUT: - return "TIMEOUT"; - case PACKET_QUEUE_TEST_BAD_HEADER: return "BAD_PACKET_HEADER"; - case PACKET_QUEUE_TEST_DATA_MISMATCH: - return "DATA_MISMATCH"; - default: - return "UNKNOWN"; - } -} - -inline uint64_t get_64b_result(uint32_t* buf, uint32_t index) { - return (((uint64_t)buf[index]) << 32) | buf[index+1]; -} - -inline uint64_t get_64b_result(const std::vector& vec, uint32_t index) { - return (((uint64_t)vec[index]) << 32) | vec[index+1]; -} - -#define TX_TEST_IDX_TOT_DATA_WORDS PQ_TEST_MISC_INDEX + 1 -#define TX_TEST_IDX_NPKT PQ_TEST_MISC_INDEX + 3 -#define TX_TEST_IDX_WORDS_FLUSHED PQ_TEST_MISC_INDEX + 5 -#define TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 7 -#define TX_TEST_IDX_MANY_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 9 -#define TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER PQ_TEST_MISC_INDEX + 11 -// #define TX_TEST_IDX_ PQ_TEST_MISC_INDEX + -// #define TX_TEST_IDX_ PQ_TEST_MISC_INDEX + - -enum class pkt_dest_size_choices_t { - RANDOM=0, - SAME_START_RNDROBIN_FIX_SIZE=1 // max packet size used -}; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_tx_ubench.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_tx_ubench.cpp index 0832c67a7c1..d9991ed8b67 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_tx_ubench.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_tx_ubench.cpp @@ -107,9 +107,9 @@ void kernel_main() { target_address = base_target_address; zero_l1_buf(test_results, test_results_size_bytes); - test_results[PQ_TEST_STATUS_INDEX] = PACKET_QUEUE_TEST_STARTED; - test_results[PQ_TEST_MISC_INDEX] = 0xff000000; - test_results[PQ_TEST_MISC_INDEX + 1] = 0xcc000000 | src_endpoint_id; + test_results[TT_FABRIC_STATUS_INDEX] = TT_FABRIC_STATUS_STARTED; + test_results[TT_FABRIC_MISC_INDEX] = 0xff000000; + test_results[TT_FABRIC_MISC_INDEX + 1] = 0xcc000000 | src_endpoint_id; zero_l1_buf( reinterpret_cast(data_buffer_start_addr), data_buffer_size_words * PACKET_WORD_SIZE_BYTES); @@ -199,11 +199,11 @@ void kernel_main() { uint64_t cycles_elapsed = get_timestamp() - start_timestamp; uint64_t num_packets = packet_count; - set_64b_result(test_results, data_words_sent, PQ_TEST_WORD_CNT_INDEX); - set_64b_result(test_results, cycles_elapsed, PQ_TEST_CYCLES_INDEX); + set_64b_result(test_results, data_words_sent, TT_FABRIC_WORD_CNT_INDEX); + set_64b_result(test_results, cycles_elapsed, TT_FABRIC_CYCLES_INDEX); set_64b_result(test_results, total_data_words, TX_TEST_IDX_TOT_DATA_WORDS); set_64b_result(test_results, num_packets, TX_TEST_IDX_NPKT); - test_results[PQ_TEST_STATUS_INDEX] = PACKET_QUEUE_TEST_PASS; - test_results[PQ_TEST_MISC_INDEX] = packet_count; + test_results[TT_FABRIC_STATUS_INDEX] = TT_FABRIC_STATUS_PASS; + test_results[TT_FABRIC_MISC_INDEX] = packet_count; } diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp index fa061868bca..f055d0a9833 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp @@ -6,6 +6,7 @@ #include #include +#include "hw/inc/tt_fabric_status.h" #include "llrt.hpp" static inline std::string to_string(pkt_dest_size_choices_t choice) { @@ -25,3 +26,11 @@ static inline void log_phys_coord_to_json(nlohmann::json& config, const std::vec static inline void log_phys_coord_to_json(nlohmann::json& config, const CoreCoord& phys_core, const std::string& name) { config[name] = fmt::format("({}, {})", phys_core.x, phys_core.y); } + +inline uint64_t get_64b_result(uint32_t* buf, uint32_t index) { + return (((uint64_t)buf[index]) << 32) | buf[index+1]; +} + +inline uint64_t get_64b_result(const std::vector& vec, uint32_t index) { + return (((uint64_t)vec[index]) << 32) | vec[index+1]; +} diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux.cpp index 03f804ce55f..05a35add66a 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux.cpp @@ -8,8 +8,7 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" #include "llrt.hpp" using std::vector; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux_2level.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux_2level.cpp index 63105c881cc..dc4a8f132fd 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux_2level.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux_2level.cpp @@ -7,7 +7,7 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" +#include "test_common.hpp" #include "llrt.hpp" using std::vector; 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..8ac6dbd69b3 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 @@ -9,8 +9,8 @@ #include "tt_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 "tt_fabric/hw/inc/tt_fabric_status.h" +#include "test_common.hpp" #include "eth_l1_address_map.h" #include "tt_fabric/hw/inc/tt_fabric_interface.h" @@ -542,12 +542,8 @@ int main(int argc, char** argv) { for (uint32_t i = 0; i < num_src_endpoints; i++) { tx_results.push_back(tt::llrt::read_hex_vec_from_core( device_map[test_device_id_l]->id(), tx_phys_core[i], test_results_addr, 128)); - log_info( - LogTest, - "TX{} status = {}", - i, - packet_queue_test_status_to_string(tx_results[i][PQ_TEST_STATUS_INDEX])); - pass &= (tx_results[i][PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + log_info(LogTest, "TX{} status = {}", i, tt_fabric_status_to_string(tx_results[i][TT_FABRIC_STATUS_INDEX])); + pass &= (tx_results[i][TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); } /* TODO: Need to add these once control plane api is available to @@ -556,15 +552,15 @@ int main(int argc, char** argv) { tt::llrt::read_hex_vec_from_core( device_map[test_device_id_l]->id(), tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "L Router status = {}", - packet_queue_test_status_to_string(router_results[PQ_TEST_STATUS_INDEX])); pass &= - (router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); vector r_router_results = tt::llrt::read_hex_vec_from_core( device_map[test_device_id_r]->id(), r_tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "R Router status = {}", - packet_queue_test_status_to_string(r_router_results[PQ_TEST_STATUS_INDEX])); pass &= - (r_router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(r_router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (r_router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); */ for (auto active_device : device_map) { pass &= tt_metal::CloseDevice(active_device.second); @@ -575,12 +571,12 @@ int main(int argc, char** argv) { uint64_t total_tx_words_sent = 0; uint64_t total_rx_words_checked = 0; for (uint32_t i = 0; i < num_src_endpoints; i++) { - uint64_t tx_words_sent = get_64b_result(tx_results[i], PQ_TEST_WORD_CNT_INDEX); + uint64_t tx_words_sent = get_64b_result(tx_results[i], TT_FABRIC_WORD_CNT_INDEX); total_tx_words_sent += tx_words_sent; - uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], PQ_TEST_CYCLES_INDEX); + uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], TT_FABRIC_CYCLES_INDEX); double tx_bw = ((double)tx_words_sent) * PACKET_WORD_SIZE_BYTES / tx_elapsed_cycles; total_tx_bw += tx_bw; - uint64_t iter = get_64b_result(tx_results[i], PQ_TEST_ITER_INDEX); + uint64_t iter = get_64b_result(tx_results[i], TT_FABRIC_ITER_INDEX); // uint64_t zero_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER); // uint64_t few_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER); // uint64_t many_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_MANY_DATA_WORDS_SENT_ITER); 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 052f8b39ed8..a0e91bd4dc2 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 @@ -10,8 +10,8 @@ #include "tt_fabric/mesh_graph.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 "tt_fabric/hw/inc/tt_fabric_status.h" +#include "test_common.hpp" #include "eth_l1_address_map.h" #include "tt_fabric/hw/inc/tt_fabric_interface.h" #include @@ -869,7 +869,8 @@ typedef struct test_traffic { num_cores_to_skip = (num_rx_workers + num_links_to_use - 1) / num_links_to_use; } // Assumes uniform worker grid across receiver chips - rx_workers = rx_devices[0]->select_worker_cores(dest_routers, num_links_to_use, num_rx_workers, num_cores_to_skip); + rx_workers = + rx_devices[0]->select_worker_cores(dest_routers, num_links_to_use, num_rx_workers, num_cores_to_skip); // TODO: not the most optimum selection, might impact somewhat in bidirectional mode controller_logical_core = tx_device->select_random_worker_cores(1)[0]; @@ -1085,8 +1086,8 @@ typedef struct test_traffic { tx_device->physical_chip_id, (uint32_t)tx_device->logical_chip_id, i, - packet_queue_test_status_to_string(tx_results[i][PQ_TEST_STATUS_INDEX])); - pass &= (tx_results[i][PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(tx_results[i][TT_FABRIC_STATUS_INDEX])); + pass &= (tx_results[i][TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); } // collect rx results @@ -1101,8 +1102,8 @@ typedef struct test_traffic { rx_devices[d]->physical_chip_id, (uint32_t)rx_devices[d]->logical_chip_id, i, - packet_queue_test_status_to_string(rx_results[d][i][PQ_TEST_STATUS_INDEX])); - pass &= (rx_results[d][i][PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(rx_results[d][i][TT_FABRIC_STATUS_INDEX])); + pass &= (rx_results[d][i][TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); } } @@ -1120,10 +1121,10 @@ typedef struct test_traffic { num_tx_packets = 0; for (auto j : rx_to_tx_map[i]) { - num_tx_words += get_64b_result(tx_results[j], PQ_TEST_WORD_CNT_INDEX); + num_tx_words += get_64b_result(tx_results[j], TT_FABRIC_WORD_CNT_INDEX); num_tx_packets += get_64b_result(tx_results[j], TX_TEST_IDX_NPKT); } - pass &= (get_64b_result(rx_results[d][i], PQ_TEST_WORD_CNT_INDEX) == num_tx_words); + pass &= (get_64b_result(rx_results[d][i], TT_FABRIC_WORD_CNT_INDEX) == num_tx_words); pass &= (get_64b_result(rx_results[d][i], TX_TEST_IDX_NPKT) == num_tx_packets); if (!pass) { @@ -1142,12 +1143,12 @@ typedef struct test_traffic { uint64_t total_rx_words_checked = 0; uint64_t max_tx_elapsed_cycles = 0; for (uint32_t i = 0; i < num_tx_workers; i++) { - uint64_t tx_words_sent = get_64b_result(tx_results[i], PQ_TEST_WORD_CNT_INDEX); + uint64_t tx_words_sent = get_64b_result(tx_results[i], TT_FABRIC_WORD_CNT_INDEX); total_tx_words_sent += tx_words_sent; - uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], PQ_TEST_CYCLES_INDEX); + uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], TT_FABRIC_CYCLES_INDEX); double tx_bw = ((double)tx_words_sent) * PACKET_WORD_SIZE_BYTES / tx_elapsed_cycles; total_tx_bw += tx_bw; - uint64_t iter = get_64b_result(tx_results[i], PQ_TEST_ITER_INDEX); + uint64_t iter = get_64b_result(tx_results[i], TT_FABRIC_ITER_INDEX); max_tx_elapsed_cycles = std::max(max_tx_elapsed_cycles, tx_elapsed_cycles); // uint64_t zero_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER); // uint64_t few_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER); @@ -1182,7 +1183,7 @@ typedef struct test_traffic { total_tx_bw_2 = ((double)total_tx_words_sent) * PACKET_WORD_SIZE_BYTES / max_tx_elapsed_cycles; for (uint32_t d = 0; d < rx_devices.size(); d++) { for (uint32_t i = 0; i < num_rx_workers; i++) { - uint64_t words_received = get_64b_result(rx_results[d][i], PQ_TEST_WORD_CNT_INDEX); + uint64_t words_received = get_64b_result(rx_results[d][i], TT_FABRIC_WORD_CNT_INDEX); uint32_t num_tx = rx_to_tx_map[i].size(); log_info( LogTest, @@ -1761,15 +1762,15 @@ int main(int argc, char **argv) { tt::llrt::read_hex_vec_from_core( device->id(), tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "L Router status = {}", - packet_queue_test_status_to_string(router_results[PQ_TEST_STATUS_INDEX])); pass &= - (router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); vector r_router_results = tt::llrt::read_hex_vec_from_core( device_r->id(), r_tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "R Router status = {}", - packet_queue_test_status_to_string(r_router_results[PQ_TEST_STATUS_INDEX])); pass &= - (r_router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(r_router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (r_router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); */ // close devices 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 14425045b9f..cf140eeaf80 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 @@ -9,8 +9,8 @@ #include "tt_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 "tt_fabric/hw/inc/tt_fabric_status.h" +#include "test_common.hpp" #include "eth_l1_address_map.h" #include "tt_fabric/hw/inc/tt_fabric_interface.h" @@ -577,12 +577,8 @@ int main(int argc, char** argv) { for (uint32_t i = 0; i < num_src_endpoints; i++) { tx_results.push_back(tt::llrt::read_hex_vec_from_core( device_map[test_device_id_l]->id(), tx_phys_core[i], test_results_addr, 128)); - log_info( - LogTest, - "TX{} status = {}", - i, - packet_queue_test_status_to_string(tx_results[i][PQ_TEST_STATUS_INDEX])); - pass &= (tx_results[i][PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + log_info(LogTest, "TX{} status = {}", i, tt_fabric_status_to_string(tx_results[i][TT_FABRIC_STATUS_INDEX])); + pass &= (tx_results[i][TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); } /* TODO: Need to add these once control plane api is available to @@ -591,15 +587,15 @@ int main(int argc, char** argv) { tt::llrt::read_hex_vec_from_core( device_map[test_device_id_l]->id(), tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "L Router status = {}", - packet_queue_test_status_to_string(router_results[PQ_TEST_STATUS_INDEX])); pass &= - (router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); vector r_router_results = tt::llrt::read_hex_vec_from_core( device_map[test_device_id_r]->id(), r_tunneler_phys_core, tunneler_test_results_addr, 128); log_info(LogTest, "R Router status = {}", - packet_queue_test_status_to_string(r_router_results[PQ_TEST_STATUS_INDEX])); pass &= - (r_router_results[PQ_TEST_STATUS_INDEX] == PACKET_QUEUE_TEST_PASS); + tt_fabric_status_to_string(r_router_results[TT_FABRIC_STATUS_INDEX])); pass &= + (r_router_results[TT_FABRIC_STATUS_INDEX] == TT_FABRIC_STATUS_PASS); */ for (auto active_device : device_map) { pass &= tt_metal::CloseDevice(active_device.second); @@ -610,12 +606,12 @@ int main(int argc, char** argv) { uint64_t total_tx_words_sent = 0; uint64_t total_rx_words_checked = 0; for (uint32_t i = 0; i < num_src_endpoints; i++) { - uint64_t tx_words_sent = get_64b_result(tx_results[i], PQ_TEST_WORD_CNT_INDEX); + uint64_t tx_words_sent = get_64b_result(tx_results[i], TT_FABRIC_WORD_CNT_INDEX); total_tx_words_sent += tx_words_sent; - uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], PQ_TEST_CYCLES_INDEX); + uint64_t tx_elapsed_cycles = get_64b_result(tx_results[i], TT_FABRIC_CYCLES_INDEX); double tx_bw = ((double)tx_words_sent) * PACKET_WORD_SIZE_BYTES / tx_elapsed_cycles; total_tx_bw += tx_bw; - uint64_t iter = get_64b_result(tx_results[i], PQ_TEST_ITER_INDEX); + uint64_t iter = get_64b_result(tx_results[i], TT_FABRIC_ITER_INDEX); // uint64_t zero_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER); // uint64_t few_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER); // uint64_t many_data_sent_iter = get_64b_result(tx_results[i], TX_TEST_IDX_MANY_DATA_WORDS_SENT_ITER); diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tx_rx.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tx_rx.cpp index e0e200af967..a645b972fa6 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tx_rx.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tx_rx.cpp @@ -7,7 +7,7 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" +#include "test_common.hpp" #include "utils.hpp" #include "llrt.hpp" diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_2ep.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_2ep.cpp index f96ca0c8528..99d271f3ce0 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_2ep.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_2ep.cpp @@ -7,9 +7,8 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" #include -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" using std::vector; using namespace tt; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_4ep.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_4ep.cpp index c1945c1b5aa..8c70290d9c3 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_4ep.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_bi_tunnel_4ep.cpp @@ -7,9 +7,8 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" #include -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" using std::vector; using namespace tt; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_loopback_tunnel.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_loopback_tunnel.cpp index 9348333bd56..0b9cf4ae5b4 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_loopback_tunnel.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_loopback_tunnel.cpp @@ -7,9 +7,8 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" #include -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" using std::vector; using namespace tt; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_mux_demux.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_mux_demux.cpp index cf6fb4609e6..11eda9992de 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_mux_demux.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_mux_demux.cpp @@ -8,8 +8,7 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" using std::vector; using namespace tt; diff --git a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_uni_tunnel.cpp b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_uni_tunnel.cpp index a837a0be959..32d69fb8586 100644 --- a/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_uni_tunnel.cpp +++ b/tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_vc_uni_tunnel.cpp @@ -8,8 +8,7 @@ #include #include #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp" -#include "kernels/traffic_gen_test.hpp" -#include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp" +#include "test_common.hpp" using std::vector; using namespace tt; diff --git a/tt_fabric/hw/inc/tt_fabric_status.h b/tt_fabric/hw/inc/tt_fabric_status.h new file mode 100644 index 00000000000..5f415112755 --- /dev/null +++ b/tt_fabric/hw/inc/tt_fabric_status.h @@ -0,0 +1,45 @@ +// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc. +// +// SPDX-License-Identifier: Apache-2.0 + +#pragma once +#include + +constexpr uint32_t TT_FABRIC_STAUS_MASK = 0xabc00000; +constexpr uint32_t TT_FABRIC_STATUS_STARTED = TT_FABRIC_STAUS_MASK | 0x0; +constexpr uint32_t TT_FABRIC_STATUS_PASS = TT_FABRIC_STAUS_MASK | 0x1; +constexpr uint32_t TT_FABRIC_STATUS_TIMEOUT = TT_FABRIC_STAUS_MASK | 0xdead0; +constexpr uint32_t TT_FABRIC_STATUS_BAD_HEADER = TT_FABRIC_STAUS_MASK | 0xdead1; +constexpr uint32_t TT_FABRIC_STATUS_DATA_MISMATCH = TT_FABRIC_STAUS_MASK | 0x3; + +// indexes of return values in test results buffer +constexpr uint32_t TT_FABRIC_STATUS_INDEX = 0; +constexpr uint32_t TT_FABRIC_WORD_CNT_INDEX = 2; +constexpr uint32_t TT_FABRIC_CYCLES_INDEX = 4; +constexpr uint32_t TT_FABRIC_ITER_INDEX = 6; +constexpr uint32_t TT_FABRIC_MISC_INDEX = 16; + +inline std::string_view tt_fabric_status_to_string(uint32_t status) { + switch (status) { + case TT_FABRIC_STATUS_STARTED: return "STARTED"; + case TT_FABRIC_STATUS_PASS: return "DONE/OK"; + case TT_FABRIC_STATUS_TIMEOUT: return "TIMEOUT"; + case TT_FABRIC_STATUS_BAD_HEADER: return "BAD_PACKET_HEADER"; + case TT_FABRIC_STATUS_DATA_MISMATCH: return "DATA_MISMATCH"; + default: return "UNKNOWN"; + } +} + +constexpr uint32_t TX_TEST_IDX_TOT_DATA_WORDS = TT_FABRIC_MISC_INDEX + 1; +constexpr uint32_t TX_TEST_IDX_NPKT = TT_FABRIC_MISC_INDEX + 3; +constexpr uint32_t TX_TEST_IDX_WORDS_FLUSHED = TT_FABRIC_MISC_INDEX + 5; +constexpr uint32_t TX_TEST_IDX_FEW_DATA_WORDS_SENT_ITER = TT_FABRIC_MISC_INDEX + 7; +constexpr uint32_t TX_TEST_IDX_MANY_DATA_WORDS_SENT_ITER = TT_FABRIC_MISC_INDEX + 9; +constexpr uint32_t TX_TEST_IDX_ZERO_DATA_WORDS_SENT_ITER = TT_FABRIC_MISC_INDEX + 11; +// constexpr uint32_t TX_TEST_IDX_ = TT_FABRIC_MISC_INDEX + ; +// constexpr uint32_t TX_TEST_IDX_ = TT_FABRIC_MISC_INDEX + ; + +enum class pkt_dest_size_choices_t { + RANDOM = 0, + SAME_START_RNDROBIN_FIX_SIZE = 1 // max packet size used +}; diff --git a/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp b/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp index 31c75c4329b..c211c6f0133 100644 --- a/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp +++ b/tt_fabric/impl/kernels/tt_fabric_gatekeeper.cpp @@ -5,6 +5,7 @@ // clang-format off #include "dataflow_api.h" #include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_fabric/hw/inc/tt_fabric_status.h" #include "debug/dprint.h" // clang-format on @@ -19,20 +20,6 @@ constexpr uint32_t timeout_cycles = get_compile_time_arg_val(5); uint32_t sync_val; uint32_t router_mask; -constexpr uint32_t PACKET_QUEUE_STAUS_MASK = 0xabc00000; -constexpr uint32_t PACKET_QUEUE_TEST_STARTED = PACKET_QUEUE_STAUS_MASK | 0x0; -constexpr uint32_t PACKET_QUEUE_TEST_PASS = PACKET_QUEUE_STAUS_MASK | 0x1; -constexpr uint32_t PACKET_QUEUE_TEST_TIMEOUT = PACKET_QUEUE_STAUS_MASK | 0xdead0; -constexpr uint32_t PACKET_QUEUE_TEST_BAD_HEADER = PACKET_QUEUE_STAUS_MASK | 0xdead1; -constexpr uint32_t PACKET_QUEUE_TEST_DATA_MISMATCH = PACKET_QUEUE_STAUS_MASK | 0x3; - -// indexes of return values in test results buffer -constexpr uint32_t PQ_TEST_STATUS_INDEX = 0; -constexpr uint32_t PQ_TEST_WORD_CNT_INDEX = 2; -constexpr uint32_t PQ_TEST_CYCLES_INDEX = 4; -constexpr uint32_t PQ_TEST_ITER_INDEX = 6; -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); volatile tt_l1_ptr fabric_router_l1_config_t* routing_table = @@ -436,11 +423,11 @@ void kernel_main() { tt_fabric_init(); - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_STARTED); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000000); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 1, 0xbb000000); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 2, 0xAABBCCDD); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 3, 0xDDCCBBAA); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_STARTED); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000000); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 1, 0xbb000000); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 2, 0xAABBCCDD); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 3, 0xDDCCBBAA); zero_l1_buf((tt_l1_ptr uint32_t*)&gk_info->gk_msg_buf, FVCC_BUF_SIZE_BYTES); zero_l1_buf((tt_l1_ptr uint32_t*)socket_info, sizeof(socket_info_t)); @@ -477,7 +464,7 @@ void kernel_main() { gk_msg_buf_advance_rdptr((ctrl_chan_msg_buf*)msg_buf); loop_count = 0; } else { - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_BAD_HEADER); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_BAD_HEADER); return; } } @@ -498,11 +485,11 @@ void kernel_main() { DPRINT << "Gatekeeper messages processed " << total_messages_procesed << ENDL(); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000002); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000002); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000003); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000003); - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_PASS); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_PASS); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff00005); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff00005); } diff --git a/tt_fabric/impl/kernels/tt_fabric_router.cpp b/tt_fabric/impl/kernels/tt_fabric_router.cpp index 5453c5f6ca3..0eeb7879f9d 100644 --- a/tt_fabric/impl/kernels/tt_fabric_router.cpp +++ b/tt_fabric/impl/kernels/tt_fabric_router.cpp @@ -5,6 +5,7 @@ // clang-format off #include "dataflow_api.h" #include "tt_fabric/hw/inc/tt_fabric.h" +#include "tt_fabric/hw/inc/tt_fabric_status.h" // clang-format on using namespace tt::tt_fabric; @@ -28,20 +29,6 @@ uint32_t router_mask; uint32_t gk_message_addr_l; uint32_t gk_message_addr_h; -constexpr uint32_t PACKET_QUEUE_STAUS_MASK = 0xabc00000; -constexpr uint32_t PACKET_QUEUE_TEST_STARTED = PACKET_QUEUE_STAUS_MASK | 0x0; -constexpr uint32_t PACKET_QUEUE_TEST_PASS = PACKET_QUEUE_STAUS_MASK | 0x1; -constexpr uint32_t PACKET_QUEUE_TEST_TIMEOUT = PACKET_QUEUE_STAUS_MASK | 0xdead0; -constexpr uint32_t PACKET_QUEUE_TEST_BAD_HEADER = PACKET_QUEUE_STAUS_MASK | 0xdead1; -constexpr uint32_t PACKET_QUEUE_TEST_DATA_MISMATCH = PACKET_QUEUE_STAUS_MASK | 0x3; - -// indexes of return values in test results buffer -constexpr uint32_t PQ_TEST_STATUS_INDEX = 0; -constexpr uint32_t PQ_TEST_WORD_CNT_INDEX = 2; -constexpr uint32_t PQ_TEST_CYCLES_INDEX = 4; -constexpr uint32_t PQ_TEST_ITER_INDEX = 6; -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 = @@ -90,11 +77,11 @@ void kernel_main() { tt_fabric_init(); - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_STARTED); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000000); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 1, 0xbb000000); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 2, 0xAABBCCDD); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX + 3, 0xDDCCBBAA); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_STARTED); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000000); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 1, 0xbb000000); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 2, 0xAABBCCDD); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX + 3, 0xDDCCBBAA); router_state.sync_in = 0; router_state.sync_out = 0; @@ -102,9 +89,9 @@ void kernel_main() { zero_l1_buf((tt_l1_ptr uint32_t*)fvc_consumer_req_buf, sizeof(chan_req_buf)); zero_l1_buf((tt_l1_ptr uint32_t*)FVCC_IN_BUF_START, FVCC_IN_BUF_SIZE); zero_l1_buf((tt_l1_ptr uint32_t*)FVCC_OUT_BUF_START, FVCC_OUT_BUF_SIZE); - write_kernel_status(kernel_status, PQ_TEST_WORD_CNT_INDEX, (uint32_t)&router_state); - write_kernel_status(kernel_status, PQ_TEST_WORD_CNT_INDEX + 1, (uint32_t)&fvc_consumer_state); - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX + 1, (uint32_t)&fvc_producer_state); + write_kernel_status(kernel_status, TT_FABRIC_WORD_CNT_INDEX, (uint32_t)&router_state); + write_kernel_status(kernel_status, TT_FABRIC_WORD_CNT_INDEX + 1, (uint32_t)&fvc_consumer_state); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX + 1, (uint32_t)&fvc_producer_state); fvc_consumer_state.init(FABRIC_ROUTER_DATA_BUF_START, fvc_data_buf_size_words / 2); fvc_producer_state.init( @@ -121,14 +108,14 @@ void kernel_main() { #endif if (!wait_all_src_dest_ready(&router_state, timeout_cycles)) { - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_TIMEOUT); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_TIMEOUT); return; } notify_gatekeeper(); uint64_t start_timestamp = get_timestamp(); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000001); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000001); uint32_t loop_count = 0; uint32_t launch_msg_rd_ptr = *GET_MAILBOX_ADDRESS_DEV(launch_msg_rd_ptr); @@ -172,7 +159,7 @@ void kernel_main() { fvc_producer_state.process_inbound_packet(); loop_count = 0; } else if (fvc_producer_state.packet_corrupted) { - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_BAD_HEADER); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_BAD_HEADER); return; } @@ -200,16 +187,16 @@ void kernel_main() { } uint64_t cycles_elapsed = fvc_producer_state.packet_timestamp - start_timestamp; - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000002); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000002); - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff000003); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff000003); - set_64b_result(kernel_status, cycles_elapsed, PQ_TEST_CYCLES_INDEX); + set_64b_result(kernel_status, cycles_elapsed, TT_FABRIC_CYCLES_INDEX); if (fvc_consumer_state.packet_in_progress) { - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_TIMEOUT); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_TIMEOUT); } else { - write_kernel_status(kernel_status, PQ_TEST_STATUS_INDEX, PACKET_QUEUE_TEST_PASS); + write_kernel_status(kernel_status, TT_FABRIC_STATUS_INDEX, TT_FABRIC_STATUS_PASS); } - write_kernel_status(kernel_status, PQ_TEST_MISC_INDEX, 0xff00005); + write_kernel_status(kernel_status, TT_FABRIC_MISC_INDEX, 0xff00005); } diff --git a/tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp b/tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp index d86086ad78d..f7be23a8d36 100644 --- a/tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp +++ b/tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp @@ -3,6 +3,7 @@ // SPDX-License-Identifier: Apache-2.0 #pragma once +#include constexpr uint32_t PACKET_WORD_SIZE_BYTES = 16; constexpr uint32_t MAX_SWITCH_FAN_IN = 4; @@ -32,6 +33,16 @@ constexpr uint32_t PQ_TEST_ITER_INDEX = 6; constexpr uint32_t PQ_TEST_MISC_INDEX = 16; +inline std::string_view packet_queue_test_status_to_string(uint32_t status) { + switch (status) { + case PACKET_QUEUE_TEST_STARTED: return "STARTED"; + case PACKET_QUEUE_TEST_PASS: return "DONE/OK"; + case PACKET_QUEUE_TEST_TIMEOUT: return "TIMEOUT"; + case PACKET_QUEUE_TEST_DATA_MISMATCH: return "DATA_MISMATCH"; + default: return "UNKNOWN"; + } +} + enum DispatchPacketFlag : uint32_t { PACKET_CMD_START = (0x1 << 1), PACKET_CMD_END = (0x1 << 2),