Skip to content

Commit

Permalink
#0: Minor code review fixes.
Browse files Browse the repository at this point in the history
  • Loading branch information
@spoojaryTT
spoojaryTT committed Nov 15, 2024
1 parent 31c7de2 commit a1c797a
Show file tree
Hide file tree
Showing 7 changed files with 30 additions and 42 deletions.
10 changes: 1 addition & 9 deletions tt_metal/hw/firmware/src/brisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -347,15 +347,7 @@ int main() {
int32_t num_words = ((uint)__ldm_data_end - (uint)__ldm_data_start) >> 2;
l1_to_local_mem_copy((uint*)__ldm_data_start, (uint tt_l1_ptr*)MEM_BRISC_INIT_LOCAL_L1_BASE_SCRATCH, num_words);

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;
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);
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);
noc_bank_table_init();

mailboxes->launch_msg_rd_ptr = 0; // Initialize the rdptr to 0
noc_index = 0;
Expand Down
10 changes: 1 addition & 9 deletions tt_metal/hw/firmware/src/erisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,15 +46,7 @@ void __attribute__((noinline)) Application(void) {
// TODO: need to find free space that routing FW is not using
wzerorange(__ldm_bss_start, __ldm_bss_end);

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;
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);
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);
noc_bank_table_init();

risc_init();
noc_init(MEM_NOC_ATOMIC_RET_VAL_ADDR);
Expand Down
10 changes: 1 addition & 9 deletions tt_metal/hw/firmware/src/idle_erisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -110,15 +110,7 @@ int main() {
local_mem_ptr[i] = l1_data_ptr[i];
}

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;
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);
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);
noc_bank_table_init();

risc_init();

Expand Down
10 changes: 1 addition & 9 deletions tt_metal/hw/firmware/src/ncrisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -83,15 +83,7 @@ int main(int argc, char *argv[]) {
int32_t num_words = ((uint)__ldm_data_end - (uint)__ldm_data_start) >> 2;
l1_to_local_mem_copy((uint *)__ldm_data_start, (uint tt_l1_ptr *)MEM_NCRISC_INIT_LOCAL_L1_BASE_SCRATCH, num_words);

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;
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);
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);
noc_bank_table_init();

risc_init();

Expand Down
20 changes: 20 additions & 0 deletions tt_metal/hw/inc/firmware_common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#include "dev_mem_map.h"
#include "hostdevcommon/kernel_structs.h"
#include "dev_msgs.h"
#include "noc/noc_parameters.h"

extern uint32_t __ldm_bss_start[];
extern uint32_t __ldm_bss_end[];
Expand All @@ -21,6 +22,11 @@ extern uint32_t __ldm_data_end[];
extern void (* __init_array_start[])();
extern void (* __init_array_end[])();

extern uint16_t dram_bank_to_noc_xy[NUM_NOCS][NUM_DRAM_BANKS];
extern int32_t bank_to_dram_offset[NUM_DRAM_BANKS];
extern uint16_t l1_bank_to_noc_xy[NUM_NOCS][NUM_L1_BANKS];
extern int32_t bank_to_l1_offset[NUM_L1_BANKS];

extern void kernel_init(uint32_t kernel_init);
extern void kernel_launch(uint32_t kernel_base_addr);

Expand Down Expand Up @@ -61,6 +67,20 @@ inline void firmware_kernel_common_init(void *init_local_l1_base) {
(**fptr)();
}
}

void noc_bank_table_init()
{
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;
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);
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);
}

FORCE_INLINE
uint32_t firmware_config_init(tt_l1_ptr mailboxes_t* const mailboxes, uint32_t core_type_index, uint32_t dispatch_class) {

Expand Down
8 changes: 4 additions & 4 deletions tt_metal/impl/device/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -401,11 +401,11 @@ void Device::build_firmware() {
log_debug(tt::LogMetal, "Building base firmware for device {}", this->id_);
ZoneScoped;

this->generate_global_arrays();
this->generate_device_bank_to_noc_tables();
jit_build_set(this->firmware_build_states_, nullptr);
}

void Device::initialize_global_array(const HalProgrammableCoreType &core_type, CoreCoord phys_core)
void Device::initialize_device_bank_to_noc_tables(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 l1_noc_addr = MEM_BANK_TO_NOC_XY_SCRATCH + (dram_bank_to_noc_xy_.size() * sizeof(uint16_t));
Expand All @@ -419,7 +419,7 @@ void Device::initialize_global_array(const HalProgrammableCoreType &core_type, C
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);
this->initialize_device_bank_to_noc_tables(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);

Expand Down Expand Up @@ -3555,7 +3555,7 @@ void Device::MarkAllocationsSafe() {
tt::tt_metal::allocator::mark_allocations_safe(*this->get_initialized_allocator());
}

void Device::generate_global_arrays()
void Device::generate_device_bank_to_noc_tables()
{
const size_t num_dram_banks = this->num_banks(BufferType::DRAM);
std::vector<CoreCoord> dram_noc_coord_per_bank(num_dram_banks);
Expand Down
4 changes: 2 additions & 2 deletions tt_metal/impl/device/device.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -231,7 +231,7 @@ class Device {
// machine inf
float sfpu_inf() const;

void generate_global_arrays();
void generate_device_bank_to_noc_tables();
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 Expand Up @@ -259,7 +259,7 @@ class Device {
void initialize_build();
void initialize_device_kernel_defines();
void build_firmware();
void initialize_global_array(const HalProgrammableCoreType &core_type, CoreCoord phys_core);
void initialize_device_bank_to_noc_tables(const HalProgrammableCoreType &core_type, CoreCoord phys_core);
void initialize_firmware(const HalProgrammableCoreType &core_type, CoreCoord phys_core, launch_msg_t *launch_msg, go_msg_t* go_msg);
void reset_cores();
void initialize_and_launch_firmware();
Expand Down

0 comments on commit a1c797a

Please sign in to comment.