From a13b98d4a82aea2da5bf7f883d12214f866d6e39 Mon Sep 17 00:00:00 2001 From: David Ma Date: Mon, 16 Dec 2024 23:33:44 +0000 Subject: [PATCH] #0: Revert some unnecessary parity hacks --- .../tt_metal/common/command_queue_fixture.hpp | 18 +--------- tt_metal/impl/device/device.cpp | 1 - tt_metal/impl/dispatch/command_queue.cpp | 36 +++++++++---------- tt_metal/impl/dispatch/command_queue.hpp | 3 -- 4 files changed, 17 insertions(+), 41 deletions(-) diff --git a/tests/tt_metal/tt_metal/common/command_queue_fixture.hpp b/tests/tt_metal/tt_metal/common/command_queue_fixture.hpp index c2928506aa4f..efec2c625342 100644 --- a/tests/tt_metal/tt_metal/common/command_queue_fixture.hpp +++ b/tests/tt_metal/tt_metal/common/command_queue_fixture.hpp @@ -72,14 +72,6 @@ class CommandQueueSingleCardFixture : virtual public DispatchFixture { this->validate_dispatch_mode(); this->arch_ = tt::get_arch_from_string(tt::test_utils::get_umd_arch_name()); this->create_devices(); - // Temp workaround until switch is flipped - if (!slow_dispatch_) { - for (Device* device : devices_) { - for (int cq_idx = 0; cq_idx < device->num_hw_cqs(); cq_idx++) { - device->hw_command_queue(cq_idx); - } - } - } } void TearDown() override { tt::tt_metal::detail::CloseDevices(reserved_devices_); } @@ -90,7 +82,7 @@ class CommandQueueSingleCardFixture : virtual public DispatchFixture { if (slow_dispatch) { tt::log_info( tt::LogTest, "This suite can only be run with fast dispatch or TT_METAL_SLOW_DISPATCH_MODE unset"); - this->slow_dispatch_ = true; + this->slow_dispatch_ = false; GTEST_SKIP(); } } @@ -122,14 +114,6 @@ class CommandQueueSingleCardTraceFixture : virtual public CommandQueueSingleCard this->validate_dispatch_mode(); this->arch_ = tt::get_arch_from_string(tt::test_utils::get_umd_arch_name()); this->create_devices(90000000); - // Temp workaround until switch is flipped - if (!slow_dispatch_) { - for (Device* device : devices_) { - for (int cq_idx = 0; cq_idx < device->num_hw_cqs(); cq_idx++) { - device->hw_command_queue(cq_idx); - } - } - } } }; diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index a687b55b58e7..103a63e053f0 100644 --- a/tt_metal/impl/device/device.cpp +++ b/tt_metal/impl/device/device.cpp @@ -3654,7 +3654,6 @@ HWCommandQueue& Device::hw_command_queue(size_t cq_id) { detail::DispatchStateCheck(true); TT_FATAL( cq_id < hw_command_queues_.size(), "cq_id {} is out of range", cq_id ); TT_FATAL(this->is_initialized(), "Device has not been initialized, did you forget to call InitializeDevice?"); - hw_command_queues_[cq_id]->update_dispatch_core(); return *hw_command_queues_[cq_id]; } diff --git a/tt_metal/impl/dispatch/command_queue.cpp b/tt_metal/impl/dispatch/command_queue.cpp index 5cdc0a875095..afc6b344740b 100644 --- a/tt_metal/impl/dispatch/command_queue.cpp +++ b/tt_metal/impl/dispatch/command_queue.cpp @@ -1962,6 +1962,22 @@ HWCommandQueue::HWCommandQueue(Device* device, uint32_t id, NOC noc_index) : this->size_B = this->size_B / 4; } + CoreCoord enqueue_program_dispatch_core; + CoreType core_type = dispatch_core_manager::instance().get_dispatch_core_type(device->id()); + if (this->device->num_hw_cqs() == 1 or core_type == CoreType::WORKER) { + // dispatch_s exists with this configuration. Workers write to dispatch_s + enqueue_program_dispatch_core = dispatch_core_manager::instance().dispatcher_s_core(device->id(), channel, id); + } + else { + if (device->is_mmio_capable()) { + enqueue_program_dispatch_core = dispatch_core_manager::instance().dispatcher_core(device->id(), channel, id); + } else { + enqueue_program_dispatch_core = dispatch_core_manager::instance().dispatcher_d_core(device->id(), channel, id); + } + } + this->virtual_enqueue_program_dispatch_core = + device->virtual_core_from_logical_core(enqueue_program_dispatch_core, core_type); + tt_cxy_pair completion_q_writer_location = dispatch_core_manager::instance().completion_queue_writer_core(device->id(), channel, this->id); @@ -2005,26 +2021,6 @@ void HWCommandQueue::set_go_signal_noc_data_on_dispatch(const vector_memcpy_alig this->manager.fetch_queue_write(cmd_sequence_sizeB, this->id); } -void HWCommandQueue::update_dispatch_core() { - CoreCoord enqueue_program_dispatch_core; - CoreType core_type = dispatch_core_manager::instance().get_dispatch_core_type(device->id()); - uint16_t channel = tt::Cluster::instance().get_assigned_channel_for_device(device->id()); - if (this->device->num_hw_cqs() == 1 or core_type == CoreType::WORKER) { - // dispatch_s exists with this configuration. Workers write to dispatch_s - enqueue_program_dispatch_core = dispatch_core_manager::instance().dispatcher_s_core(device->id(), channel, id); - } else { - if (device->is_mmio_capable()) { - enqueue_program_dispatch_core = - dispatch_core_manager::instance().dispatcher_core(device->id(), channel, id); - } else { - enqueue_program_dispatch_core = - dispatch_core_manager::instance().dispatcher_d_core(device->id(), channel, id); - } - } - this->virtual_enqueue_program_dispatch_core = - device->virtual_core_from_logical_core(enqueue_program_dispatch_core, core_type); -} - void HWCommandQueue::reset_worker_state(bool reset_launch_msg_state) { auto num_sub_devices = device->num_sub_devices(); uint32_t go_signals_cmd_size = 0; diff --git a/tt_metal/impl/dispatch/command_queue.hpp b/tt_metal/impl/dispatch/command_queue.hpp index 9a9fa5036f3f..661e3d6d4992 100644 --- a/tt_metal/impl/dispatch/command_queue.hpp +++ b/tt_metal/impl/dispatch/command_queue.hpp @@ -528,9 +528,6 @@ class HWCommandQueue { void set_go_signal_noc_data_on_dispatch(const vector_memcpy_aligned& go_signal_noc_data); void reset_worker_state(bool reset_launch_msg_state); - // Temporary to get dispatch assignments matching with new implementation - void update_dispatch_core(); - private: uint32_t id; uint32_t size_B;