#16339: parameterize dispatch_constants

- Add DispatchSettings

- Add DispatchConstants. Struct with only static constexpr
  values. Later to remove those values from dispatch_constants
  which is doing memory/address layout calculations

- Put align methods into helper file

- Some namespace pollution reduced

tests/tt_metal/tt_metal/common/command_queue_fixture.hpp

@@ -15,6 +15,7 @@
 #include "tt_metal/impl/kernels/kernel.hpp"
 #include "tt_metal/common/tt_backend_api_types.hpp"
 #include "tt_metal/llrt/rtoptions.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 class CommandQueueFixture : public DispatchFixture {
 protected:

tests/tt_metal/tt_metal/common/device_fixture.hpp

@@ -11,6 +11,7 @@
 #include "tt_metal/detail/tt_metal.hpp"
 #include "tt_metal/test_utils/env_vars.hpp"
 #include "tt_metal/impl/device/device_pool.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 class DeviceFixture : public DispatchFixture {
 protected:

tests/tt_metal/tt_metal/common/dispatch_fixture.hpp

@@ -12,6 +12,7 @@
 #include "tt_metal/impl/dispatch/command_queue.hpp"
 #include "tt_metal/impl/device/device.hpp"
 #include "tt_metal/impl/device/device_pool.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 // A dispatch-agnostic test fixture
 class DispatchFixture : public ::testing::Test {

tests/tt_metal/tt_metal/common/matmul_test_utils.hpp

@@ -10,6 +10,7 @@
 #include "tt_metal/common/test_tiles.hpp"
 #include "hostdevcommon/common_values.hpp"
 #include "tt_metal/impl/dispatch/command_queue.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using namespace tt;
 

tests/tt_metal/tt_metal/common/multi_device_fixture.hpp

@@ -11,6 +11,7 @@
 #include "umd/device/types/cluster_descriptor_types.h"
 #include "tt_metal/test_utils/env_vars.hpp"
 #include "tt_metal/impl/device/device_pool.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 class MultiDeviceFixture : public DispatchFixture {
 protected:

tests/tt_metal/tt_metal/debug_tools/debug_tools_test_utils.hpp

@@ -5,6 +5,7 @@
 #pragma once
 
 #include "host_api.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 // Helper function to open a file as an fstream, and check that it was opened properly.
 inline bool OpenFile(string &file_name, std::fstream &file_stream, std::ios_base::openmode mode) {

tests/tt_metal/tt_metal/dispatch/CMakeLists.txt

@@ -2,13 +2,15 @@ add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/dispatch_buffer)
 add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/dispatch_event)
 add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/dispatch_program)
 add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/dispatch_trace)
+add_subdirectory(${CMAKE_CURRENT_SOURCE_DIR}/dispatch_util)
 
 add_executable(
     unit_tests_dispatch
     $<TARGET_OBJECTS:unit_tests_dispatch_buffer_o>
     $<TARGET_OBJECTS:unit_tests_dispatch_event_o>
     $<TARGET_OBJECTS:unit_tests_dispatch_program_o>
     $<TARGET_OBJECTS:unit_tests_dispatch_trace_o>
+    $<TARGET_OBJECTS:unit_tests_dispatch_util_o>
 )
 TT_ENABLE_UNITY_BUILD(unit_tests_dispatch)
 

tests/tt_metal/tt_metal/dispatch/dispatch_util/CMakeLists.txt

@@ -0,0 +1,31 @@
+set(UNIT_TESTS_DISPATCH_UTIL_SRC ${CMAKE_CURRENT_SOURCE_DIR}/test_dispatch_settings.cpp)
+
+add_library(unit_tests_dispatch_util_o STATIC ${UNIT_TESTS_DISPATCH_UTIL_SRC})
+
+target_link_libraries(unit_tests_dispatch_util_o PRIVATE test_metal_common_libs)
+
+target_include_directories(
+    unit_tests_dispatch_util_o
+    PRIVATE
+        ${PROJECT_SOURCE_DIR}
+        ${PROJECT_SOURCE_DIR}/tt_metal
+        ${PROJECT_SOURCE_DIR}/tests
+        ${PROJECT_SOURCE_DIR}/tests/tt_metal/tt_metal/common
+        ${PROJECT_SOURCE_DIR}/tests/tt_metal/tt_metal/dispatch
+        ${PROJECT_SOURCE_DIR}/tests/tt_metal/tt_metal/dispatch/common
+        ${CMAKE_CURRENT_SOURCE_DIR}
+        ${CMAKE_CURRENT_SOURCE_DIR}/common
+)
+
+add_executable(unit_tests_dispatch_util $<TARGET_OBJECTS:unit_tests_dispatch_util_o>)
+
+target_link_libraries(unit_tests_dispatch_util PRIVATE test_metal_common_libs)
+
+set_target_properties(
+    unit_tests_dispatch_util
+    PROPERTIES
+        RUNTIME_OUTPUT_DIRECTORY
+            ${PROJECT_BINARY_DIR}/test/tt_metal
+)
+
+TT_ENABLE_UNITY_BUILD(unit_tests_dispatch_util)

tests/tt_metal/tt_metal/dispatch/dispatch_util/test_dispatch_settings.cpp

@@ -0,0 +1,134 @@
+// SPDX-FileCopyrightText: © 2024 Tenstorrent Inc.
+//
+// SPDX-License-Identifier: Apache-2.0
+
+#include <array>
+#include <stdexcept>
+#include "command_queue_fixture.hpp"
+#include "common/logger.hpp"
+#include "dispatch/dispatch_constants.hpp"
+#include "gtest/gtest.h"
+#include "llrt/hal.hpp"
+#include "tt_metal/impl/dispatch/util/include/dispatch_settings.hpp"
+#include "tt_metal/impl/dispatch/command_queue_interface.hpp"
+#include "umd/device/tt_core_coordinates.h"
+
+using namespace tt::tt_metal::dispatch;
+
+// Loop through test_func for WORKER, ETH X 1, 2 CQs
+void ForEachCoreTypeXHWCQs(const std::function<void(const CoreType& core_type, const uint32_t num_hw_cqs)>& test_func) {
+    static constexpr auto core_types_to_test = std::array<CoreType, 2>{CoreType::WORKER, CoreType::ETH};
+    static constexpr auto num_hw_cqs_to_test = std::array<uint32_t, 2>{1, 2};
+
+    for (const auto& core_type : core_types_to_test) {
+        if (core_type == CoreType::ETH &&
+            hal.get_programmable_core_type_index(tt::tt_metal::HalProgrammableCoreType::IDLE_ETH) == -1) {
+            // This device does not have the eth core
+            tt::log_info(tt::LogTest, "IDLE_ETH core type is not on this device");
+            continue;
+        }
+        for (const auto& num_hw_cqs : num_hw_cqs_to_test) {
+            test_func(core_type, num_hw_cqs);
+        }
+    }
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsDefaultParity) {
+    ForEachCoreTypeXHWCQs([&](const CoreType& core_type, uint32_t num_hw_cqs) {
+        auto settings = DispatchSettings::defaults(core_type, tt::Cluster::instance(), num_hw_cqs);
+
+        const auto& old_constants = dispatch_constants::get(core_type, num_hw_cqs);
+
+        ASSERT_EQ(settings.num_hw_cqs_, num_hw_cqs);
+
+        ASSERT_EQ(settings.prefetch_q_entries_, old_constants.prefetch_q_entries());
+        ASSERT_EQ(settings.prefetch_q_size_, old_constants.prefetch_q_size());
+        ASSERT_EQ(settings.prefetch_max_cmd_size_, old_constants.max_prefetch_command_size());
+        ASSERT_EQ(settings.prefetch_cmddat_q_size_, old_constants.cmddat_q_size());
+        ASSERT_EQ(settings.prefetch_scratch_db_size_, old_constants.scratch_db_size());
+
+        ASSERT_EQ(settings.prefetch_d_buffer_size_, old_constants.prefetch_d_buffer_size());
+        ASSERT_EQ(settings.prefetch_d_pages_, old_constants.prefetch_d_buffer_pages());
+        ASSERT_EQ(settings.prefetch_d_blocks_, dispatch_constants::PREFETCH_D_BUFFER_BLOCKS);
+
+        ASSERT_EQ(settings.tunneling_buffer_size_ / num_hw_cqs, old_constants.mux_buffer_size(num_hw_cqs));
+        ASSERT_EQ(settings.tunneling_buffer_pages_ / num_hw_cqs, old_constants.mux_buffer_pages(num_hw_cqs));
+
+        ASSERT_EQ(settings.dispatch_pages_, old_constants.dispatch_buffer_pages());
+        ASSERT_EQ(settings.dispatch_pages_per_block_, dispatch_constants::DISPATCH_BUFFER_SIZE_BLOCKS);
+
+        ASSERT_EQ(settings.dispatch_s_buffer_size_, old_constants.dispatch_s_buffer_size());
+        ASSERT_EQ(settings.dispatch_s_buffer_pages_, old_constants.dispatch_s_buffer_pages());
+
+        EXPECT_NO_THROW(settings.build());
+    });
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsDefaultUnsupportedCoreType) {
+    const auto unsupported_core = CoreType::ARC;
+    EXPECT_THROW(DispatchSettings::defaults(unsupported_core, tt::Cluster::instance(), 1), std::runtime_error);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsMissingArgs) {
+    DispatchSettings settings;
+    EXPECT_THROW(settings.build(), std::runtime_error);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsEq) {
+    static constexpr uint32_t hw_cqs = 2;
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    auto settings_2 = settings; // Copy
+    EXPECT_EQ(settings, settings_2);
+    settings_2.dispatch_size_ += 1;
+    EXPECT_NE(settings, settings_2);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsSetPrefetchDBuffer) {
+    static constexpr uint32_t hw_cqs = 2;
+    static constexpr uint32_t expected_buffer_bytes = 0xcafe;
+    static constexpr uint32_t expected_page_count = expected_buffer_bytes / (1 << DispatchConstants::PREFETCH_D_BUFFER_LOG_PAGE_SIZE);
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    settings.prefetch_d_buffer_size(expected_buffer_bytes);
+    EXPECT_EQ(settings.prefetch_d_buffer_size_, expected_buffer_bytes);
+    EXPECT_EQ(settings.prefetch_d_pages_, expected_page_count);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsSetPrefetchQBuffer) {
+    static constexpr uint32_t hw_cqs = 2;
+    static constexpr uint32_t expected_buffer_entries = 0x1000;
+    static constexpr uint32_t expected_buffer_bytes = expected_buffer_entries * sizeof(DispatchConstants::prefetch_q_entry_type);
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    settings.prefetch_q_entries(expected_buffer_entries);
+    EXPECT_EQ(settings.prefetch_q_entries_, expected_buffer_entries);
+    EXPECT_EQ(settings.prefetch_q_size_, expected_buffer_bytes);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsSetDispatchBuffer) {
+    static constexpr uint32_t hw_cqs = 2;
+    static constexpr uint32_t expected_buffer_bytes = 0x2000;
+    static constexpr uint32_t expected_page_count = expected_buffer_bytes / (1 << DispatchConstants::DISPATCH_BUFFER_LOG_PAGE_SIZE);
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    settings.dispatch_size(expected_buffer_bytes);
+    EXPECT_EQ(settings.dispatch_size_, expected_buffer_bytes);
+    EXPECT_EQ(settings.dispatch_pages_, expected_page_count);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsSetDispatchSBuffer) {
+    static constexpr uint32_t hw_cqs = 2;
+    static constexpr uint32_t expected_buffer_bytes = 0x2000;
+    static constexpr uint32_t expected_page_count = expected_buffer_bytes / (1 << DispatchConstants::DISPATCH_S_BUFFER_LOG_PAGE_SIZE);
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    settings.dispatch_s_buffer_size(expected_buffer_bytes);
+    EXPECT_EQ(settings.dispatch_s_buffer_size_, expected_buffer_bytes);
+    EXPECT_EQ(settings.dispatch_s_buffer_pages_, expected_page_count);
+}
+
+TEST_F(CommandQueueSingleCardFixture, TestDispatchSettingsSetTunnelerBuffer) {
+    static constexpr uint32_t hw_cqs = 2;
+    static constexpr uint32_t expected_buffer_bytes = 0x2000;
+    static constexpr uint32_t expected_page_count = expected_buffer_bytes / (1 << DispatchConstants::PREFETCH_D_BUFFER_LOG_PAGE_SIZE);
+    auto settings = DispatchSettings::worker_defaults(tt::Cluster::instance(), hw_cqs);
+    settings.tunneling_buffer_size(expected_buffer_bytes);
+    EXPECT_EQ(settings.tunneling_buffer_size_, expected_buffer_bytes);
+    EXPECT_EQ(settings.tunneling_buffer_pages_, expected_page_count);
+}

tests/tt_metal/tt_metal/eth/test_basic_eth.cpp

@@ -12,6 +12,7 @@
 #include "host_api.hpp"
 #include "tt_metal/impl/kernels/kernel.hpp"
 #include "tt_metal/test_utils/stimulus.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 // TODO: ARCH_NAME specific, must remove
 #include "eth_l1_address_map.h"

tests/tt_metal/tt_metal/perf_microbenchmark/common/util.hpp

@@ -12,6 +12,7 @@
 #include "tt_metal/host_api.hpp"
 #include "tt_metal/hostdevcommon/dprint_common.h"
 #include "llrt/hal.hpp"
+#include "llrt/llrt.hpp"
 
 inline uint64_t get_t0_to_any_riscfw_end_cycle(tt::tt_metal::Device* device, const tt::tt_metal::Program& program) {
 #if defined(TRACY_ENABLE)

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h

@@ -14,6 +14,7 @@
 #include "noc/noc_parameters.h"
 
 #include "tt_metal/llrt/hal.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 extern bool debug_g;
 extern bool use_coherent_data_g;

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/test_prefetcher.cpp

@@ -17,6 +17,7 @@
 #include "tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/traffic_gen_test.hpp"
 
 #include "llrt/hal.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 #define CQ_PREFETCH_CMD_BARE_MIN_SIZE tt::tt_metal::hal.get_alignment(tt::tt_metal::HalMemType::HOST)
 

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_bi_tunnel.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
 #include "tt_metal/impl/device/device.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_common.hpp

@@ -6,6 +6,7 @@
 
 #include <nlohmann/json.hpp>
 #include "tt_metal/common/core_coord.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 static inline std::string to_string(pkt_dest_size_choices_t choice) {
     switch (choice) {

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux.cpp

@@ -10,6 +10,7 @@
 #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 "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_mux_demux_2level.cpp

@@ -8,6 +8,7 @@
 #include "tt_metal/llrt/rtoptions.hpp"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tunnel_1cq.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tunnel_2cq.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/device/device.hpp"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_tx_rx.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
 #include "utils.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_uni_tunnel.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
 #include "tt_metal/impl/dispatch/kernels/packet_queue_ctrl.hpp"
 #include "kernels/traffic_gen_test.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using std::vector;
 using namespace tt;

tests/tt_metal/tt_metal/perf_microbenchmark/routing/test_uni_tunnel_single_chip.cpp

@@ -9,6 +9,7 @@
 #include "tt_metal/impl/dispatch/cq_commands.hpp"
 #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"
 
 using std::vector;
 using namespace tt;

tt_metal/impl/CMakeLists.txt

@@ -30,6 +30,7 @@ set(IMPL_SRC
     ${CMAKE_CURRENT_SOURCE_DIR}/dispatch/kernel_config/demux.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/dispatch/kernel_config/eth_router.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/dispatch/kernel_config/eth_tunneler.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/dispatch/util/dispatch_settings.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/debug/dprint_server.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/debug/noc_logging.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/debug/watcher_server.cpp

tt_metal/impl/debug/noc_logging.cpp

@@ -13,6 +13,7 @@
 #include "debug_helpers.hpp"
 #include "hostdevcommon/dprint_common.h"
 #include "tt_metal/impl/device/device.hpp"
+#include "tt_metal/llrt/llrt.hpp"
 
 using namespace tt::tt_metal;