diff --git a/tt_metal/impl/device/device.cpp b/tt_metal/impl/device/device.cpp index 4881b75f9323..1361518cd4b6 100644 --- 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 index 3bb4fd1e6b49..c97ea165a38b 100644 --- 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& dram_bank_offset_map, std::vector& l1_bank_map, std::vector& l1_bank_offset_map, - uint32_t allocator_alignment) { + uint32_t allocator_alignment, + const std::vector& 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& dram_bank_offset_map, std::vector& l1_bank_map, std::vector& l1_bank_offset_map, - uint32_t allocator_alignment) { + uint32_t allocator_alignment, + const std::vector& 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 index 4dee07a44ab3..646b9bc36171 100644 --- 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& dram_bank_offset_map, std::vector& l1_bank_map, std::vector& l1_bank_offset_map, - uint32_t allocator_alignment); + uint32_t allocator_alignment, + const std::vector& eth_chan_map); void jit_build_genfiles_descriptors(const JitBuildEnv& env, JitBuildOptions& options);