#3764: opID update

tenstorrent · Jul 11, 2024 · 0cead39 · 0cead39
1 parent 4035964
commit 0cead39
Show file tree

Hide file tree

Showing 3 changed files with 3 additions and 2 deletions.
diff --git a/tt_metal/hw/firmware/src/brisc.cc b/tt_metal/hw/firmware/src/brisc.cc
@@ -368,7 +368,7 @@ int main() {
 
         {
             DeviceZoneScopedMainN("BRISC-FW");
-            DeviceZoneSetCounter(mailboxes->launch.host_assigned_op_id);
+            DeviceZoneSetCounter(mailboxes->launch.kernel_config.host_assigned_op_id);
 
             // Copies from L1 to IRAM on chips where NCRISC has IRAM
             l1_to_ncrisc_iram_copy(mailboxes->launch.kernel_config.ncrisc_kernel_size16, ncrisc_kernel_start_offset16);

diff --git a/tt_metal/hw/inc/dev_msgs.h b/tt_metal/hw/inc/dev_msgs.h
@@ -94,6 +94,7 @@ struct kernel_config_msg_t {
     volatile uint8_t dispatch_core_y;
     volatile uint8_t exit_erisc_kernel;
     volatile uint8_t pad1;
+    volatile uint16_t pad2;
 } __attribute__((packed));
 
 struct go_msg_t {

diff --git a/tt_metal/tt_metal.cpp b/tt_metal/tt_metal.cpp
@@ -532,7 +532,7 @@ void LaunchProgram(Device *device, Program &program, bool wait_until_cores_done)
         for (const auto &[core_type, logical_cores] : logical_cores_used_in_program) {
             for (const auto &logical_core : logical_cores) {
                 launch_msg_t *msg = &program.kernels_on_core(logical_core, core_type)->launch_msg;
-                msg->host_assigned_op_id = program.get_global_id();
+                msg->kernel_config.host_assigned_op_id = program.get_global_id();
                 auto physical_core = device->physical_core_from_logical_core(logical_core, core_type);
                 not_done_cores.insert(physical_core);
                 tt::llrt::write_launch_msg_to_core(device->id(), physical_core, msg);