From 63e04c135564922cd7c1978e132de6b5602457ef Mon Sep 17 00:00:00 2001 From: Bryan Wilder Field Lozano Date: Thu, 7 Nov 2024 22:37:20 -0800 Subject: [PATCH] #14608: l1_address_params init to use Hal (#14872) --- tt_metal/llrt/blackhole/bh_hal_active_eth.cpp | 8 ++++++-- tt_metal/llrt/hal.hpp | 1 + tt_metal/llrt/tt_cluster.cpp | 18 ++++++------------ tt_metal/llrt/tt_cluster.hpp | 7 +------ tt_metal/llrt/wormhole/wh_hal_active_eth.cpp | 8 ++++++-- 5 files changed, 20 insertions(+), 22 deletions(-) diff --git a/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp b/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp index 763e6567b01..7e17ccfaca0 100644 --- a/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp +++ b/tt_metal/llrt/blackhole/bh_hal_active_eth.cpp @@ -6,6 +6,8 @@ #define COMPILE_FOR_IDLE_ERISC +#include + #include "llrt/hal.hpp" #include "llrt/blackhole/bh_hal.hpp" #include "hw/inc/blackhole/core_config.h" @@ -29,7 +31,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); - mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch); mem_map_bases[utils::underlying_type(HalL1MemAddrType::WATCHER)] = GET_ETH_MAILBOX_ADDRESS_HOST(watcher); @@ -40,10 +42,11 @@ HalCoreInfoType create_active_eth_mem_map() { mem_map_bases[utils::underlying_type(HalL1MemAddrType::CORE_INFO)] = GET_ETH_MAILBOX_ADDRESS_HOST(core_info); mem_map_bases[utils::underlying_type(HalL1MemAddrType::GO_MSG)] = GET_ETH_MAILBOX_ADDRESS_HOST(go_message); mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch_msg_rd_ptr); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::FW_VERSION_ADDR)] = eth_l1_mem::address_map::FW_VERSION_ADDR; std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); - mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::WATCHER)] = sizeof(watcher_msg_t); @@ -53,6 +56,7 @@ HalCoreInfoType create_active_eth_mem_map() { mem_map_sizes[utils::underlying_type(HalL1MemAddrType::UNRESERVED)] = eth_l1_mem::address_map::MAX_SIZE - eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::GO_MSG)] = sizeof(go_msg_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = sizeof(uint32_t); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::FW_VERSION_ADDR)] = sizeof(std::uint32_t); // TODO (abhullar): This should be NumEthDispatchClasses std::vector> processor_classes(1); diff --git a/tt_metal/llrt/hal.hpp b/tt_metal/llrt/hal.hpp index 875252ca8da..1ba7a104e84 100644 --- a/tt_metal/llrt/hal.hpp +++ b/tt_metal/llrt/hal.hpp @@ -50,6 +50,7 @@ enum class HalL1MemAddrType : uint8_t { CORE_INFO, GO_MSG, LAUNCH_MSG_BUFFER_RD_PTR, + FW_VERSION_ADDR, // Really only applicable to active eth core right now COUNT // Keep this last so it always indicates number of enum options }; diff --git a/tt_metal/llrt/tt_cluster.cpp b/tt_metal/llrt/tt_cluster.cpp index 03ed414f7a2..715ab3c974e 100644 --- a/tt_metal/llrt/tt_cluster.cpp +++ b/tt_metal/llrt/tt_cluster.cpp @@ -280,6 +280,12 @@ void Cluster::open_driver(const bool &skip_driver_allocs) { } std::uint32_t dram_barrier_base = tt_metal::hal.get_dev_addr(tt_metal::HalDramMemAddrType::DRAM_BARRIER); device_driver->set_device_dram_address_params(tt_device_dram_address_params{dram_barrier_base}); + + l1_address_params.tensix_l1_barrier_base = tt_metal::hal.get_dev_addr(tt_metal::HalProgrammableCoreType::TENSIX, tt_metal::HalL1MemAddrType::BARRIER); + if (tt_metal::hal.get_arch() != tt::ARCH::GRAYSKULL) { + l1_address_params.eth_l1_barrier_base = tt_metal::hal.get_dev_addr(tt_metal::HalProgrammableCoreType::ACTIVE_ETH, tt_metal::HalL1MemAddrType::BARRIER); + l1_address_params.fw_version_addr = tt_metal::hal.get_dev_addr(tt_metal::HalProgrammableCoreType::ACTIVE_ETH, tt_metal::HalL1MemAddrType::FW_VERSION_ADDR); + } device_driver->set_device_l1_address_params(l1_address_params); this->get_metal_desc_from_tt_desc( @@ -343,18 +349,6 @@ uint32_t Cluster::get_harvested_rows(chip_id_t chip) const { } } -void Cluster::verify_eth_fw() const { - for (const auto &[chip, mmio_device_id] : this->device_to_mmio_device_) { - std::vector fw_versions; - for (const CoreCoord ð_core : get_soc_desc(chip).ethernet_cores) { - uint32_t val; - read_core(&val, sizeof(uint32_t), tt_cxy_pair(chip, eth_core), eth_l1_mem::address_map::FW_VERSION_ADDR); - fw_versions.push_back(val); - } - verify_sw_fw_versions(chip, SW_VERSION, fw_versions); - } -} - int Cluster::get_device_aiclk(const chip_id_t &chip_id) const { if (this->arch_ == tt::ARCH::BLACKHOLE) { // For Blackhole bring up remove AICLK query due to lack of ARC message support diff --git a/tt_metal/llrt/tt_cluster.hpp b/tt_metal/llrt/tt_cluster.hpp index 59633f28bee..f59bf5d42e9 100644 --- a/tt_metal/llrt/tt_cluster.hpp +++ b/tt_metal/llrt/tt_cluster.hpp @@ -71,7 +71,6 @@ class Cluster { } //! device driver and misc apis - void verify_eth_fw() const; void verify_sw_fw_versions(int device_id, std::uint32_t sw_version, std::vector &fw_versions) const; void deassert_risc_reset_at_core(const tt_cxy_pair &physical_chip_coord) const; @@ -284,11 +283,7 @@ class Cluster { // Mapping of each devices' ethernet routing mode std::unordered_map> device_eth_routing_info_; - tt_device_l1_address_params l1_address_params = { - (uint32_t)MEM_L1_BARRIER, - (uint32_t)eth_l1_mem::address_map::ERISC_BARRIER_BASE, - (uint32_t)eth_l1_mem::address_map::FW_VERSION_ADDR, - }; + tt_device_l1_address_params l1_address_params; std::unordered_map>> ethernet_sockets_; }; diff --git a/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp b/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp index 423098ba44f..ecb4a74a1ef 100644 --- a/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp +++ b/tt_metal/llrt/wormhole/wh_hal_active_eth.cpp @@ -6,6 +6,8 @@ #define COMPILE_FOR_ERISC +#include + #include "llrt/hal.hpp" #include "llrt/wormhole/wh_hal.hpp" #include "hw/inc/wormhole/core_config.h" @@ -29,7 +31,7 @@ HalCoreInfoType create_active_eth_mem_map() { std::vector mem_map_bases; mem_map_bases.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); - mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = MEM_L1_BARRIER; + mem_map_bases[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_BASE; mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch); mem_map_bases[utils::underlying_type(HalL1MemAddrType::WATCHER)] = GET_ETH_MAILBOX_ADDRESS_HOST(watcher); @@ -40,10 +42,11 @@ HalCoreInfoType create_active_eth_mem_map() { mem_map_bases[utils::underlying_type(HalL1MemAddrType::CORE_INFO)] = GET_ETH_MAILBOX_ADDRESS_HOST(core_info); mem_map_bases[utils::underlying_type(HalL1MemAddrType::GO_MSG)] = GET_ETH_MAILBOX_ADDRESS_HOST(go_message); mem_map_bases[utils::underlying_type(HalL1MemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = GET_ETH_MAILBOX_ADDRESS_HOST(launch_msg_rd_ptr); + mem_map_bases[utils::underlying_type(HalL1MemAddrType::FW_VERSION_ADDR)] = eth_l1_mem::address_map::FW_VERSION_ADDR; std::vector mem_map_sizes; mem_map_sizes.resize(utils::underlying_type(HalL1MemAddrType::COUNT)); - mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = sizeof(uint32_t); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::BARRIER)] = eth_l1_mem::address_map::ERISC_BARRIER_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::MAILBOX)] = eth_l1_mem::address_map::ERISC_MEM_MAILBOX_SIZE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH)] = sizeof(launch_msg_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::WATCHER)] = sizeof(watcher_msg_t); @@ -53,6 +56,7 @@ HalCoreInfoType create_active_eth_mem_map() { mem_map_sizes[utils::underlying_type(HalL1MemAddrType::UNRESERVED)] = eth_l1_mem::address_map::MAX_SIZE - eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE; mem_map_sizes[utils::underlying_type(HalL1MemAddrType::GO_MSG)] = sizeof(go_msg_t); mem_map_sizes[utils::underlying_type(HalL1MemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = sizeof(uint32_t); + mem_map_sizes[utils::underlying_type(HalL1MemAddrType::FW_VERSION_ADDR)] = sizeof(std::uint32_t); std::vector> processor_classes(NumEthDispatchClasses); std::vector processor_types{0};