Skip to content

Commit

Permalink
#12844: Remove array generation in genfiles, add asserts, code cleanup
Browse files Browse the repository at this point in the history
  • Loading branch information
@spoojaryTT
spoojaryTT committed Nov 14, 2024
1 parent 6fb325a commit 5611f40
Show file tree
Hide file tree
Showing 9 changed files with 23 additions and 218 deletions.
33 changes: 0 additions & 33 deletions tt_metal/hw/firmware/src/brisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,6 @@ uint32_t tt_l1_ptr *rta_l1_base __attribute__((used));
uint32_t tt_l1_ptr *crta_l1_base __attribute__((used));
uint32_t tt_l1_ptr *sem_l1_base[ProgrammableCoreType::COUNT] __attribute__((used));

//DRAM and L1 bank offsets and noc coordinates.
uint16_t dram_bank_to_noc_xy[NUM_NOCS][NUM_DRAM_BANKS] __attribute__((used));
uint16_t l1_bank_to_noc_xy[NUM_NOCS][NUM_L1_BANKS] __attribute__((used));
int32_t bank_to_dram_offset[NUM_DRAM_BANKS] __attribute__((used));
Expand Down Expand Up @@ -350,45 +349,13 @@ int main() {

int32_t num_dram_to_noc_words = (NUM_NOCS * NUM_DRAM_BANKS) << 1;
l1_to_local_mem_copy((uint*)dram_bank_to_noc_xy, (uint tt_l1_ptr*)MEM_BANK_TO_NOC_XY_SCRATCH, num_dram_to_noc_words);

int32_t num_l1_to_noc_words = (NUM_NOCS * NUM_L1_BANKS) << 1;
//DPRINT << " base addr = " << MEM_BANK_TO_NOC_XY_SCRATCH << " num_dram_to_noc_words = " << num_dram_to_noc_words << " num_l1_to_noc_words = " << num_l1_to_noc_words << ENDL();
l1_to_local_mem_copy((uint*)l1_bank_to_noc_xy, (uint tt_l1_ptr*)(MEM_BANK_TO_NOC_XY_SCRATCH + num_dram_to_noc_words), num_l1_to_noc_words);

int32_t num_dram_offset_words = NUM_DRAM_BANKS << 2;
l1_to_local_mem_copy((uint*)bank_to_dram_offset, (uint tt_l1_ptr*)(MEM_BANK_OFFSET_SCRATCH), num_dram_offset_words);

//DPRINT << "DRAM noc_xy[0][0] = " << dram_bank_to_noc_xy[0][0] << " and " << temp_dram_bank_to_noc_xy[0][0] << ENDL();
//DPRINT << "DRAM noc_xy[0][1] = " << dram_bank_to_noc_xy[0][1] << " and " << temp_dram_bank_to_noc_xy[0][1] << ENDL();
//DPRINT << "DRAM noc_xy[0][last] = " << dram_bank_to_noc_xy[0][NUM_DRAM_BANKS-1] << " and " << temp_dram_bank_to_noc_xy[0][NUM_DRAM_BANKS-1] << ENDL();
//DPRINT << "DRAM noc_xy[0][last-1] = " << dram_bank_to_noc_xy[0][NUM_DRAM_BANKS-2] << " and " << temp_dram_bank_to_noc_xy[0][NUM_DRAM_BANKS-2] << ENDL();

//DPRINT << "DRAM noc_xy[1][0] = " << dram_bank_to_noc_xy[1][0] << " and " << temp_dram_bank_to_noc_xy[1][0] << ENDL();
//DPRINT << "DRAM noc_xy[1][1] = " << dram_bank_to_noc_xy[1][1] << " and " << temp_dram_bank_to_noc_xy[1][1] << ENDL();
//DPRINT << "DRAM noc_xy[1][last] = " << dram_bank_to_noc_xy[1][NUM_DRAM_BANKS-1] << " and " << temp_dram_bank_to_noc_xy[1][NUM_DRAM_BANKS-1] << ENDL();
//DPRINT << "DRAM noc_xy[1][last-1] = " << dram_bank_to_noc_xy[1][NUM_DRAM_BANKS-2] << " and " << temp_dram_bank_to_noc_xy[1][NUM_DRAM_BANKS-2] << ENDL();

//DPRINT << "l1 noc_xy[0][0] = " << l1_bank_to_noc_xy[0][0] << " and " << temp_l1_bank_to_noc_xy[0][0] << ENDL();
//DPRINT << "l1 noc_xy[0][1] = " << l1_bank_to_noc_xy[0][1] << " and " << temp_l1_bank_to_noc_xy[0][1] << ENDL();
//DPRINT << "l1 noc_xy[0][last] = " << l1_bank_to_noc_xy[0][NUM_L1_BANKS-1] << " and " << temp_l1_bank_to_noc_xy[0][NUM_L1_BANKS-1] << ENDL();
//DPRINT << "l1 noc_xy[0][last-1] = " << l1_bank_to_noc_xy[0][NUM_L1_BANKS-2] << " and " << temp_l1_bank_to_noc_xy[0][NUM_L1_BANKS-2] << ENDL();

//DPRINT << "l1 noc_xy[1][0] = " << l1_bank_to_noc_xy[1][0] << " and " << temp_l1_bank_to_noc_xy[1][0] << ENDL();
//DPRINT << "l1 noc_xy[1][1] = " << l1_bank_to_noc_xy[1][1] << " and " << temp_l1_bank_to_noc_xy[1][1] << ENDL();
//DPRINT << "l1 noc_xy[1][last] = " << l1_bank_to_noc_xy[1][NUM_L1_BANKS-1] << " and " << temp_l1_bank_to_noc_xy[1][NUM_L1_BANKS-1] << ENDL();
//DPRINT << "l1 noc_xy[1][last-1] = " << l1_bank_to_noc_xy[1][NUM_L1_BANKS-2] << " and " << temp_l1_bank_to_noc_xy[1][NUM_L1_BANKS-2] << ENDL();

//DPRINT << "dram_offset [0] = " << bank_to_dram_offset[0] << " and " << temp_bank_to_dram_offset[0] << ENDL();
//DPRINT << "dram_offset [1] = " << bank_to_dram_offset[1] << " and " << temp_bank_to_dram_offset[1] << ENDL();
//DPRINT << "dram_offset [last] = " << bank_to_dram_offset[NUM_DRAM_BANKS-1] << " and " << temp_bank_to_dram_offset[NUM_DRAM_BANKS-1] << ENDL();
//DPRINT << "dram_offset [last-1] = " << bank_to_dram_offset[NUM_DRAM_BANKS-2] << " and " << temp_bank_to_dram_offset[NUM_DRAM_BANKS-2] << ENDL();

int32_t num_l1_offset_words = NUM_L1_BANKS << 2;
l1_to_local_mem_copy((uint*)bank_to_l1_offset, (uint tt_l1_ptr*)(MEM_BANK_OFFSET_SCRATCH + num_dram_offset_words), num_l1_offset_words);
//DPRINT << "l1_offset [0] = " << bank_to_l1_offset[0] << " and " << temp_bank_to_l1_offset[0] << ENDL();
//DPRINT << "l1_offset [1] = " << bank_to_l1_offset[1] << " and " << temp_bank_to_l1_offset[1] << ENDL();
//DPRINT << "l1_offset [last] = " << bank_to_l1_offset[NUM_L1_BANKS-1] << " and " << temp_bank_to_l1_offset[NUM_L1_BANKS-1] << ENDL();
//DPRINT << "l1_offset [last-1] = " << bank_to_l1_offset[NUM_L1_BANKS-2] << " and " << temp_bank_to_l1_offset[NUM_L1_BANKS-2] << ENDL();

mailboxes->launch_msg_rd_ptr = 0; // Initialize the rdptr to 0
noc_index = 0;
Expand Down
1 change: 0 additions & 1 deletion tt_metal/hw/firmware/src/ncrisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -38,7 +38,6 @@ uint32_t tt_l1_ptr *rta_l1_base __attribute__((used));
uint32_t tt_l1_ptr *crta_l1_base __attribute__((used));
uint32_t tt_l1_ptr *sem_l1_base[ProgrammableCoreType::COUNT] __attribute__((used));

//DRAM and L1 bank offsets and noc coordinates.
uint16_t dram_bank_to_noc_xy[NUM_NOCS][NUM_DRAM_BANKS] __attribute__((used));
int32_t bank_to_dram_offset[NUM_DRAM_BANKS] __attribute__((used));
uint16_t l1_bank_to_noc_xy[NUM_NOCS][NUM_L1_BANKS] __attribute__((used));
Expand Down
8 changes: 5 additions & 3 deletions tt_metal/hw/inc/dataflow_api.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -10,9 +10,7 @@
#include "chlkc_unpack_tile_dims.h"
#define DATA_FORMATS_DEFINED
#endif
#if __has_include("generated_bank_to_noc_coord_mapping.h")
#include "generated_bank_to_noc_coord_mapping.h"
#endif
#include <noc/noc_parameters.h>

#include <stdint.h>

Expand Down Expand Up @@ -91,6 +89,10 @@ extern CBInterface cb_interface[NUM_CIRCULAR_BUFFERS];
#define EXCLUDE_START_X_OFFSET 8
#define DYNAMIC_NOC_DIRECTION(noc, direction) (noc == 1 ? 1 - direction : direction)

static_assert(NUM_NOCS == 2);
// "Scratch" in L1 has space allocated for 256 DRAM and L1 enteries, to store offsets and NOC XY data. (MEM_BANK_TO_NOC_XY_SCRATCH and MEM_BANK_OFFSET_SCRATCH)
static_assert((NUM_DRAM_BANKS + NUM_L1_BANKS) <= 256);

FORCE_INLINE
uint32_t align(uint32_t addr, uint32_t alignment) { return ((addr - 1) | (alignment - 1)) + 1; }

Expand Down
62 changes: 17 additions & 45 deletions tt_metal/impl/device/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,6 @@
#include <chrono>
#include <type_traits>
#include "tt_metal/host_api.hpp"
#include "tt_metal/jit_build/genfiles.hpp"
#include "tt_metal/impl/device/device.hpp"
#include "tt_metal/impl/trace/trace.hpp"
#include "tt_metal/common/core_descriptor.hpp"
Expand Down Expand Up @@ -402,31 +401,30 @@ void Device::build_firmware() {
log_debug(tt::LogMetal, "Building base firmware for device {}", this->id_);
ZoneScoped;

this->generate_device_headers(this->build_env_.get_out_firmware_root_path());
this->generate_global_arrays();
jit_build_set(this->firmware_build_states_, nullptr);
}

void Device::initialize_global_array(const HalProgrammableCoreType &core_type, CoreCoord phys_core)
{
tt::Cluster::instance().write_core(&dram_bank_to_noc_xy_[0], dram_bank_to_noc_xy_.size() * sizeof(uint16_t), tt_cxy_pair(this->id(), phys_core), MEM_BANK_TO_NOC_XY_SCRATCH);
uint64_t addr = MEM_BANK_TO_NOC_XY_SCRATCH + (dram_bank_to_noc_xy_.size() * sizeof(uint16_t));
tt::Cluster::instance().write_core(&l1_bank_to_noc_xy_[0], l1_bank_to_noc_xy_.size() * sizeof(uint16_t), tt_cxy_pair(this->id(), phys_core), addr);
uint64_t l1_noc_addr = MEM_BANK_TO_NOC_XY_SCRATCH + (dram_bank_to_noc_xy_.size() * sizeof(uint16_t));
tt::Cluster::instance().write_core(&l1_bank_to_noc_xy_[0], l1_bank_to_noc_xy_.size() * sizeof(uint16_t), tt_cxy_pair(this->id(), phys_core), l1_noc_addr);

tt::Cluster::instance().write_core(&dram_bank_offset_map_[0], dram_bank_offset_map_.size() * sizeof(int32_t), tt_cxy_pair(this->id(), phys_core), MEM_BANK_OFFSET_SCRATCH);
addr = MEM_BANK_OFFSET_SCRATCH + (dram_bank_offset_map_.size() * sizeof(int32_t));
tt::Cluster::instance().write_core(&l1_bank_offset_map_[0], l1_bank_offset_map_.size() * sizeof(int32_t), tt_cxy_pair(this->id(), phys_core), addr);
uint64_t l1_offset_addr = MEM_BANK_OFFSET_SCRATCH + (dram_bank_offset_map_.size() * sizeof(int32_t));
tt::Cluster::instance().write_core(&l1_bank_offset_map_[0], l1_bank_offset_map_.size() * sizeof(int32_t), tt_cxy_pair(this->id(), phys_core), l1_offset_addr);
}

void Device::initialize_firmware(const HalProgrammableCoreType &core_type, CoreCoord phys_core, launch_msg_t *launch_msg, go_msg_t* go_msg) {
ZoneScoped;

this->initialize_global_array(core_type, phys_core);

uint32_t core_type_idx = hal.get_programmable_core_type_index(core_type);
uint32_t processor_class_count = hal.get_processor_classes_count(core_type);

switch (core_type) {
case HalProgrammableCoreType::TENSIX: {
this->initialize_global_array(core_type, phys_core);
llrt::program_risc_startup_addr(this->id(), phys_core);
for (uint32_t processor_class = 0; processor_class < processor_class_count; processor_class++) {
auto [build_idx, num_build_states] = this->build_processor_type_to_index(core_type_idx, processor_class);
Expand Down Expand Up @@ -617,7 +615,6 @@ void Device::reset_cores() {

void Device::initialize_and_launch_firmware() {
ZoneScoped;
tt::log_info(tt::LogTest, "initialize_and_launch_firmware called for ");

launch_msg_t launch_msg;
go_msg_t go_msg;
Expand Down Expand Up @@ -3558,10 +3555,9 @@ void Device::MarkAllocationsSafe() {
tt::tt_metal::allocator::mark_allocations_safe(*this->get_initialized_allocator());
}

void Device::generate_device_headers(const std::string &path)
void Device::generate_global_arrays()
{
const size_t num_dram_banks = this->num_banks(BufferType::DRAM);
const size_t num_dram_banks_pow2 = std::pow(2, std::ceil(std::log2(num_dram_banks)));
std::vector<CoreCoord> dram_noc_coord_per_bank(num_dram_banks);
dram_bank_offset_map_.clear();
dram_bank_offset_map_.resize(num_dram_banks);
Expand All @@ -3570,7 +3566,6 @@ void Device::generate_device_headers(const std::string &path)
dram_bank_offset_map_[bank_id] = this->bank_offset(BufferType::DRAM, bank_id);
}
const size_t num_l1_banks = this->num_banks(BufferType::L1);
const size_t num_l1_banks_pow2 = std::pow(2, std::ceil(std::log2(num_l1_banks)));
std::vector<CoreCoord> l1_noc_coord_per_bank(num_l1_banks);
l1_bank_offset_map_.clear();
l1_bank_offset_map_.resize(num_l1_banks);
Expand All @@ -3581,46 +3576,23 @@ void Device::generate_device_headers(const std::string &path)

const metal_SocDescriptor& soc_d = tt::Cluster::instance().get_soc_desc(this->id());

// Generate header file in proper location
tt::log_info(tt::LogTest, "Calling genfiles for path {}", path);
jit_build_genfiles_bank_to_noc_coord_descriptor (
path,
soc_d.grid_size,
dram_noc_coord_per_bank,
dram_bank_offset_map_,
l1_noc_coord_per_bank,
l1_bank_offset_map_,
this->get_allocator_alignment()
);
this->generate_mem_bank_info(
soc_d.grid_size,
dram_noc_coord_per_bank,
l1_noc_coord_per_bank);
}

void Device::generate_mem_bank_info(
tt_xy_pair grid_size,
std::vector<CoreCoord>& dram_bank_map,
std::vector<CoreCoord>& l1_bank_map)
{
tt::log_info(tt::LogTest, "generate_mem_bank_info called for device {}\n", this->id_);
dram_bank_to_noc_xy_.clear();
dram_bank_to_noc_xy_.reserve(2 * dram_bank_map.size());
for (unsigned int noc = 0; noc < 2; noc++) {
for (unsigned int bank_id = 0; bank_id < dram_bank_map.size(); bank_id++) {
uint16_t noc_x = NOC_0_X(noc, grid_size.x, dram_bank_map[bank_id].x);
uint16_t noc_y = NOC_0_Y(noc, grid_size.y, dram_bank_map[bank_id].y);
dram_bank_to_noc_xy_.reserve(NUM_NOCS * dram_noc_coord_per_bank.size());
for (unsigned int noc = 0; noc < NUM_NOCS; noc++) {
for (unsigned int bank_id = 0; bank_id < dram_noc_coord_per_bank.size(); bank_id++) {
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, dram_noc_coord_per_bank[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, dram_noc_coord_per_bank[bank_id].y);
uint16_t xy = ((noc_y << NOC_ADDR_NODE_ID_BITS) | noc_x) << NOC_COORD_REG_OFFSET;
dram_bank_to_noc_xy_.push_back(xy);
}
}

l1_bank_to_noc_xy_.clear();
l1_bank_to_noc_xy_.reserve(2 * l1_bank_map.size());
for (unsigned int noc = 0; noc < 2; noc++) {
for (unsigned int bank_id = 0; bank_id < l1_bank_map.size(); bank_id++) {
uint16_t noc_x = NOC_0_X(noc, grid_size.x, l1_bank_map[bank_id].x);
uint16_t noc_y = NOC_0_Y(noc, grid_size.y, l1_bank_map[bank_id].y);
l1_bank_to_noc_xy_.reserve(NUM_NOCS * l1_noc_coord_per_bank.size());
for (unsigned int noc = 0; noc < NUM_NOCS; noc++) {
for (unsigned int bank_id = 0; bank_id < l1_noc_coord_per_bank.size(); bank_id++) {
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, l1_noc_coord_per_bank[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, l1_noc_coord_per_bank[bank_id].y);
uint16_t xy = ((noc_y << NOC_ADDR_NODE_ID_BITS) | noc_x) << NOC_COORD_REG_OFFSET;
l1_bank_to_noc_xy_.push_back(xy);
}
Expand Down
6 changes: 1 addition & 5 deletions tt_metal/impl/device/device.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -231,11 +231,7 @@ class Device {
// machine inf
float sfpu_inf() const;

void generate_device_headers(const std::string &path);
void generate_mem_bank_info(
tt_xy_pair grid_size,
std::vector<CoreCoord>& dram_bank_map,
std::vector<CoreCoord>& l1_bank_map);
void generate_global_arrays();
const JitBuildEnv& build_env() const { return this->build_env_; }
const string build_firmware_target_path(uint32_t programmable_core, uint32_t processor_class, int i) const;
const string build_kernel_target_path(uint32_t programmable_core, uint32_t processor_class, int i, const string& kernel_name) const;
Expand Down
1 change: 0 additions & 1 deletion tt_metal/jit_build/build.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -635,7 +635,6 @@ void JitBuildState::build(const JitBuildSettings* settings) const {
std::remove(log_file.c_str());
}

tt::log_info(tt::LogTest, "Building JitBuild for {}", out_dir);
compile(log_file, out_dir, settings);
link(log_file, out_dir);
if (this->is_fw_) {
Expand Down
1 change: 0 additions & 1 deletion tt_metal/jit_build/build.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -51,7 +51,6 @@ class JitBuildEnv {
tt::ARCH get_arch() const { return arch_; }
const string& get_root_path() const { return root_; }
const string& get_out_root_path() const { return out_root_; }
const string& get_out_firmware_root_path() const { return out_firmware_root_; }
const string& get_out_kernel_root_path() const { return out_kernel_root_; }

private:
Expand Down
120 changes: 0 additions & 120 deletions tt_metal/jit_build/genfiles.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -543,124 +543,4 @@ void jit_build_genfiles_descriptors(const JitBuildEnv& env, JitBuildOptions& opt
}
}

std::string generate_bank_to_noc_coord_descriptor_string(
tt_xy_pair grid_size,
std::vector<CoreCoord>& dram_bank_map,
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) {
stringstream ss;

ss << "// SPDX-FileCopyrightText: © 2023 Tenstorrent Inc." << endl;
ss << "//" << endl;
ss << "// SPDX-License-Identifier: Apache-2.0" << endl;
ss << endl;
ss << "/*" << endl;
ss << " * This file is autogenerated by tt-metal runtime" << endl;
ss << " * DO NOT EDIT" << endl;
ss << " * This file contains values that are visible to the device compiled code." << endl;
ss << " * CAREFUL: when included in the FW_BUILD, it defines global variables." << endl;
ss << " * When included in KERNEL_BUILD, it declares global variables." << endl;
ss << " */" << endl;
ss << endl;
ss << "#pragma once" << endl;
ss << endl;
ss << "#include <noc/noc_parameters.h>" << endl;
ss << endl;

ss << "static_assert(NUM_NOCS == 2);" << endl;
ss << endl;

ss << "#ifdef KERNEL_BUILD" << endl;
ss << endl;

ss << endl;
ss << "#else // !KERNEL_BUILD (FW_BUILD)" << 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 << " {"
<< "\t// noc=" << noc << endl;
for (unsigned int bank_id = 0; bank_id < dram_bank_map.size(); bank_id++) {
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, dram_bank_map[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, dram_bank_map[bank_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 << "int32_t bank_to_dram_offset[NUM_DRAM_BANKS] __attribute__((used)) = {" << endl;
for (unsigned int bank_id = 0; bank_id < dram_bank_map.size(); bank_id++) {
ss << " " << dram_bank_offset_map[bank_id] << "," << endl;
}
ss << "};" << endl;
ss << endl;

ss << "uint16_t l1_bank_to_noc_xy[NUM_NOCS][NUM_L1_BANKS] __attribute__((used)) = {" << endl;
for (unsigned int noc = 0; noc < 2; noc++) {
ss << " {"
<< "\t// noc=" << noc << endl;
for (unsigned int bank_id = 0; bank_id < l1_bank_map.size(); bank_id++) {
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, l1_bank_map[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, l1_bank_map[bank_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 << "int32_t bank_to_l1_offset[NUM_L1_BANKS] __attribute__((used)) = {" << endl;
for (unsigned int bank_id = 0; bank_id < l1_bank_map.size(); bank_id++) {
ss << " " << l1_bank_offset_map[bank_id] << "," << endl;
}
ss << "};" << endl;
ss << endl;

ss << "#endif // FW_BUILD" << endl;

return ss.str();
}
void jit_build_genfiles_bank_to_noc_coord_descriptor(
const string& path,
tt_xy_pair grid_size,
std::vector<CoreCoord>& dram_bank_map,
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) {
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);

fs::create_directories(path + "/brisc");
ofstream file_stream_br(path + "/brisc/generated_bank_to_noc_coord_mapping.h");
file_stream_br << output_string;
file_stream_br.close();
fs::create_directories(path + "/ncrisc");
ofstream file_stream_nc(path + "/ncrisc/generated_bank_to_noc_coord_mapping.h");
file_stream_nc << output_string;
file_stream_nc.close();
fs::create_directories(path + "/erisc");
ofstream file_stream_ec(path + "/erisc/generated_bank_to_noc_coord_mapping.h");
file_stream_ec << output_string;
file_stream_ec.close();
fs::create_directories(path + "/idle_erisc");
ofstream file_stream_iec(path + "/idle_erisc/generated_bank_to_noc_coord_mapping.h");
file_stream_iec << output_string;
file_stream_iec.close();
fs::create_directories(path + "/slave_idle_erisc");
ofstream file_stream_siec(path + "/slave_idle_erisc/generated_bank_to_noc_coord_mapping.h");
file_stream_siec << output_string;
file_stream_siec.close();
}

} // namespace tt::tt_metal
Loading

0 comments on commit 5611f40

Please sign in to comment.