Add eth chan mapping to generated file

tenstorrent · Nov 26, 2024 · a6715b4 · a6715b4
1 parent 839ecd5
commit a6715b4
Show file tree

Hide file tree

Showing 3 changed files with 29 additions and 7 deletions.
diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp
@@ -323,6 +323,8 @@ void Device::initialize_device_kernel_defines()
     this->device_kernel_defines_.emplace("PCIE_NOC_Y", std::to_string(pcie_cores[0].y));
     this->device_kernel_defines_.emplace("PCIE_NOC1_X", std::to_string(tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].x)));
     this->device_kernel_defines_.emplace("PCIE_NOC1_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(NOC::NOC_1, grid_size.x, pcie_cores[0].y)));
+
+    this->device_kernel_defines_.emplace("NUM_ETH_CHANS", std::to_string(soc_d.physical_ethernet_cores.size()));
 }
 
 void Device::initialize_build() {
@@ -3568,15 +3570,15 @@ void Device::generate_device_headers(const std::string &path) const
     const metal_SocDescriptor& soc_d = tt::Cluster::instance().get_soc_desc(this->id());
 
     // Generate header file in proper location
-    jit_build_genfiles_bank_to_noc_coord_descriptor (
+    jit_build_genfiles_bank_to_noc_coord_descriptor(
         path,
         soc_d.grid_size,
         dram_noc_coord_per_bank,
         dram_offsets_per_bank,
         l1_noc_coord_per_bank,
         l1_offset_per_bank,
-        this->get_allocator_alignment()
-    );
+        this->get_allocator_alignment(),
+        soc_d.physical_ethernet_cores);
 }
 
 size_t Device::get_device_kernel_defines_hash() {

diff --git a/tt_metal/jit_build/genfiles.cpp b/tt_metal/jit_build/genfiles.cpp
@@ -549,7 +549,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
     std::vector<int32_t>& dram_bank_offset_map,
     std::vector<CoreCoord>& l1_bank_map,
     std::vector<int32_t>& l1_bank_offset_map,
-    uint32_t allocator_alignment) {
+    uint32_t allocator_alignment,
+    const std::vector<CoreCoord>& eth_chan_map) {
     stringstream ss;
 
     ss << "// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc." << endl;
@@ -574,6 +575,7 @@ std::string generate_bank_to_noc_coord_descriptor_string(
 
     ss << "#ifdef KERNEL_BUILD" << endl;
     ss << endl;
+    ss << "extern uint16_t eth_chan_to_noc_xy[NUM_NOCS][NUM_ETH_CHANS];" << endl;
     ss << "extern uint16_t dram_bank_to_noc_xy[NUM_NOCS][NUM_DRAM_BANKS];" << endl;
     ss << "extern int32_t bank_to_dram_offset[NUM_DRAM_BANKS];" << endl;
     ss << "extern uint16_t l1_bank_to_noc_xy[NUM_NOCS][NUM_L1_BANKS];" << endl;
@@ -583,6 +585,21 @@ std::string generate_bank_to_noc_coord_descriptor_string(
     ss << "#else // !KERNEL_BUILD (FW_BUILD)" << endl;
     ss << endl;
 
+    ss << "uint16_t eth_chan_to_noc_xy[NUM_NOCS][NUM_ETH_CHANS] __attribute__((used)) = {" << endl;
+    for (unsigned int noc = 0; noc < 2; noc++) {
+        ss << "    {"
+           << "\t// noc=" << noc << endl;
+        for (unsigned int chan_id = 0; chan_id < eth_chan_map.size(); chan_id++) {
+            uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, eth_chan_map[chan_id].x);
+            uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, eth_chan_map[chan_id].y);
+            ss << "        (((" << noc_y << " << NOC_ADDR_NODE_ID_BITS) | " << noc_x << ") << NOC_COORD_REG_OFFSET),"
+               << "\t// NOC_X=" << noc_x << " NOC_Y=" << noc_y << endl;
+        }
+        ss << "    }," << endl;
+    }
+    ss << "};" << endl;
+    ss << endl;
+
     ss << "uint16_t dram_bank_to_noc_xy[NUM_NOCS][NUM_DRAM_BANKS] __attribute__((used)) = {" << endl;
     for (unsigned int noc = 0; noc < 2; noc++) {
         ss << "    {"
@@ -636,14 +653,16 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor(
     std::vector<int32_t>& dram_bank_offset_map,
     std::vector<CoreCoord>& l1_bank_map,
     std::vector<int32_t>& l1_bank_offset_map,
-    uint32_t allocator_alignment) {
+    uint32_t allocator_alignment,
+    const std::vector<CoreCoord>& eth_chan_map) {
     string output_string = generate_bank_to_noc_coord_descriptor_string(
         grid_size,
         dram_bank_map,
         dram_bank_offset_map,
         l1_bank_map,
         l1_bank_offset_map,
-        allocator_alignment);
+        allocator_alignment,
+        eth_chan_map);
 
     fs::create_directories(path + "/brisc");
     ofstream file_stream_br(path + "/brisc/generated_bank_to_noc_coord_mapping.h");

diff --git a/tt_metal/jit_build/genfiles.hpp b/tt_metal/jit_build/genfiles.hpp
@@ -28,7 +28,8 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor(
     std::vector<int32_t>& dram_bank_offset_map,
     std::vector<CoreCoord>& l1_bank_map,
     std::vector<int32_t>& l1_bank_offset_map,
-    uint32_t allocator_alignment);
+    uint32_t allocator_alignment,
+    const std::vector<CoreCoord>& eth_chan_map);
 
 void jit_build_genfiles_descriptors(const JitBuildEnv& env, JitBuildOptions& options);