#0: Revert some unnecessary parity hacks

tenstorrent · Dec 16, 2024 · a13b98d · a13b98d
1 parent 0f43d2e
commit a13b98d
Show file tree

Hide file tree

Showing 4 changed files with 17 additions and 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
@@ -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
@@ -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
@@ -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
@@ -528,9 +528,6 @@ class HWCommandQueue {
     void set_go_signal_noc_data_on_dispatch(const vector_memcpy_aligned<uint32_t>& 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;