Skip to content

Commit

Permalink
Enable eth cores on BH and allow loading FW onto active eriscs
Browse files Browse the repository at this point in the history
  • Loading branch information
@abhullar-tt
abhullar-tt committed Feb 26, 2025
1 parent 6d4a570 commit 7409fd8
Show file tree
Hide file tree
Showing 19 changed files with 211 additions and 68 deletions.
30 changes: 30 additions & 0 deletions tests/tt_metal/tt_metal/common/multi_device_fixture.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,36 @@ class MultiDeviceFixture : public DispatchFixture {
void SetUp() override { this->arch_ = tt::get_arch_from_string(tt::test_utils::get_umd_arch_name()); }
};

class TwoDeviceFixture : public MultiDeviceFixture {
protected:
void SetUp() override {
this->slow_dispatch_ = true;
auto slow_dispatch = getenv("TT_METAL_SLOW_DISPATCH_MODE");
if (!slow_dispatch) {
tt::log_info(tt::LogTest, "This suite can only be run with TT_METAL_SLOW_DISPATCH_MODE set");
this->slow_dispatch_ = false;
GTEST_SKIP();
}

MultiDeviceFixture::SetUp();

const size_t num_devices = tt::tt_metal::GetNumAvailableDevices();
const size_t num_pci_devices = tt::tt_metal::GetNumPCIeDevices();
if (num_devices == 2) {
std::vector<chip_id_t> ids;
for (chip_id_t id = 0; id < num_devices; id++) {
ids.push_back(id);
}

const auto& dispatch_core_config = tt::llrt::RunTimeOptions::get_instance().get_dispatch_core_config();
tt::DevicePool::initialize(ids, 1, DEFAULT_L1_SMALL_SIZE, DEFAULT_TRACE_REGION_SIZE, dispatch_core_config);
this->devices_ = tt::DevicePool::instance().get_all_active_devices();
} else {
GTEST_SKIP() << "TwoDeviceFixture can only be run on machines with two devices";
}
}
};

class N300DeviceFixture : public MultiDeviceFixture {
protected:
void SetUp() override {
Expand Down
16 changes: 14 additions & 2 deletions tests/tt_metal/tt_metal/eth/test_buffer_movement_kernels.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -300,11 +300,14 @@ bool chip_to_chip_interleaved_buffer_transfer(

} // namespace unit_tests::erisc::kernels

TEST_F(N300DeviceFixture, ActiveEthKernelsSendDramBufferChip0ToChip1) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsSendDramBufferChip0ToChip1) {
const auto& sender_device = devices_.at(0);
const auto& receiver_device = devices_.at(1);

for (const auto& sender_eth_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_eth_core)) {
continue;
}
CoreCoord receiver_eth_core = std::get<1>(sender_device->get_connected_ethernet_core(sender_eth_core));

ASSERT_TRUE(unit_tests::erisc::kernels::chip_to_chip_dram_buffer_transfer(
Expand Down Expand Up @@ -338,11 +341,14 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsSendDramBufferChip0ToChip1) {
}
}

TEST_F(N300DeviceFixture, ActiveEthKernelsSendDramBufferChip1ToChip0) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsSendDramBufferChip1ToChip0) {
const auto& sender_device = devices_.at(1);
const auto& receiver_device = devices_.at(0);

for (const auto& sender_eth_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_eth_core)) {
continue;
}
CoreCoord receiver_eth_core = std::get<1>(sender_device->get_connected_ethernet_core(sender_eth_core));

ASSERT_TRUE(unit_tests::erisc::kernels::chip_to_chip_dram_buffer_transfer(
Expand Down Expand Up @@ -452,6 +458,9 @@ TEST_F(DeviceFixture, ActiveEthKernelsSendInterleavedBufferAllConnectedChips) {
continue;
}
for (const auto& sender_eth_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_eth_core)) {
continue;
}
auto [device_id, receiver_eth_core] = sender_device->get_connected_ethernet_core(sender_eth_core);
if (receiver_device->id() != device_id) {
continue;
Expand Down Expand Up @@ -521,6 +530,9 @@ TEST_F(CommandQueueMultiDeviceProgramFixture, ActiveEthKernelsSendDramBufferAllC
continue;
}
for (const auto& sender_eth_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_eth_core)) {
continue;
}
auto [device_id, receiver_eth_core] = sender_device->get_connected_ethernet_core(sender_eth_core);
if (receiver_device->id() != device_id) {
continue;
Expand Down
40 changes: 36 additions & 4 deletions tests/tt_metal/tt_metal/eth/test_erisc_app_direct_send.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -517,6 +517,9 @@ TEST_F(DeviceFixture, ActiveEthKernelsDirectSendAllConnectedChips) {
continue;
}
for (const auto& sender_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_core)) {
continue;
}
auto [device_id, receiver_core] = sender_device->get_connected_ethernet_core(sender_core);
if (receiver_device->id() != device_id) {
continue;
Expand Down Expand Up @@ -562,7 +565,7 @@ TEST_F(DeviceFixture, ActiveEthKernelsDirectSendAllConnectedChips) {
}
}

TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
using namespace CMAKE_UNIQUE_NAMESPACE;
const auto& device_0 = devices_.at(0);
const auto& device_1 = devices_.at(1);
Expand All @@ -571,6 +574,9 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
const size_t dst_eth_l1_byte_address = eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE;

for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
CoreCoord receiver_core = std::get<1>(device_0->get_connected_ethernet_core(sender_core));
ASSERT_TRUE(unit_tests::erisc::direct_send::eth_direct_sender_receiver_kernels(
static_cast<DispatchFixture*>(this),
Expand All @@ -592,6 +598,9 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
sender_core));
}
for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
CoreCoord receiver_core = std::get<1>(device_0->get_connected_ethernet_core(sender_core));
ASSERT_TRUE(unit_tests::erisc::direct_send::eth_direct_sender_receiver_kernels(
static_cast<DispatchFixture*>(this),
Expand All @@ -613,6 +622,9 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
sender_core));
}
for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
CoreCoord receiver_core = std::get<1>(device_0->get_connected_ethernet_core(sender_core));
ASSERT_TRUE(unit_tests::erisc::direct_send::eth_direct_sender_receiver_kernels(
static_cast<DispatchFixture*>(this),
Expand All @@ -634,6 +646,9 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
sender_core));
}
for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
CoreCoord receiver_core = std::get<1>(device_0->get_connected_ethernet_core(sender_core));
ASSERT_TRUE(unit_tests::erisc::direct_send::eth_direct_sender_receiver_kernels(
static_cast<DispatchFixture*>(this),
Expand All @@ -656,7 +671,7 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsBidirectionalDirectSend) {
}
}

TEST_F(N300DeviceFixture, ActiveEthKernelsRepeatedDirectSends) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsRepeatedDirectSends) {
using namespace CMAKE_UNIQUE_NAMESPACE;
const auto& device_0 = devices_.at(0);
const auto& device_1 = devices_.at(1);
Expand All @@ -665,6 +680,9 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsRepeatedDirectSends) {
const size_t dst_eth_l1_byte_address = eth_l1_mem::address_map::ERISC_L1_UNRESERVED_BASE;

for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
CoreCoord receiver_core = std::get<1>(device_0->get_connected_ethernet_core(sender_core));
for (int i = 0; i < 10; i++) {
ASSERT_TRUE(unit_tests::erisc::direct_send::eth_direct_sender_receiver_kernels(
Expand All @@ -691,7 +709,7 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsRepeatedDirectSends) {
}
}

TEST_F(N300DeviceFixture, ActiveEthKernelsRandomDirectSendTests) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsRandomDirectSendTests) {
using namespace CMAKE_UNIQUE_NAMESPACE;
srand(0);
const auto& device_0 = devices_.at(0);
Expand All @@ -712,6 +730,11 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsRandomDirectSendTests) {

const auto& send_chip = devices_.at(std::get<0>(it->first));
CoreCoord sender_core = std::get<1>(it->first);

if (not tt::Cluster::instance().is_ethernet_link_up(send_chip->id(), sender_core)) {
continue;
}

const auto& receiver_chip = devices_.at(std::get<0>(it->second));
CoreCoord receiver_core = std::get<1>(it->second);

Expand All @@ -736,17 +759,23 @@ TEST_F(N300DeviceFixture, ActiveEthKernelsRandomDirectSendTests) {
receiver_core));
}
}
TEST_F(N300DeviceFixture, ActiveEthKernelsRandomEthPacketSizeDirectSendTests) {
TEST_F(TwoDeviceFixture, ActiveEthKernelsRandomEthPacketSizeDirectSendTests) {
srand(0);
const auto& device_0 = devices_.at(0);
const auto& device_1 = devices_.at(1);

std::map<std::tuple<int, CoreCoord>, std::tuple<int, CoreCoord>> connectivity = {};
for (const auto& sender_core : device_0->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_0->id(), sender_core)) {
continue;
}
const auto& receiver_core = device_0->get_connected_ethernet_core(sender_core);
connectivity.insert({{0, sender_core}, receiver_core});
}
for (const auto& sender_core : device_1->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(device_1->id(), sender_core)) {
continue;
}
const auto& receiver_core = device_1->get_connected_ethernet_core(sender_core);
connectivity.insert({{1, sender_core}, receiver_core});
}
Expand Down Expand Up @@ -799,6 +828,9 @@ TEST_F(CommandQueueMultiDeviceProgramFixture, ActiveEthKernelsDirectSendAllConne
continue;
}
for (const auto& sender_core : sender_device->get_active_ethernet_cores(true)) {
if (not tt::Cluster::instance().is_ethernet_link_up(sender_device->id(), sender_core)) {
continue;
}
auto [device_id, receiver_core] = sender_device->get_connected_ethernet_core(sender_core);
if (receiver_device->id() != device_id) {
continue;
Expand Down
4 changes: 2 additions & 2 deletions tt_metal/api/tt-metalium/dev_msgs.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -308,9 +308,9 @@ struct addressable_core_t {
// This is the number of Ethernet cores on WH (Ethernet cores can be queried through Virtual Coordinates).
// All other Non Worker Cores are not accessible through virtual coordinates. Subject to change, depending on the arch.
constexpr static std::uint32_t MAX_VIRTUAL_NON_WORKER_CORES = 18;
// This is the total number of Non Worker Cores on WH (first term is Ethernet, second term is PCIe and last term is
// This is the total number of Non Worker Cores on WH (first term is DRAM, second term is PCIe and last term is
// DRAM).
constexpr static std::uint32_t MAX_NON_WORKER_CORES = MAX_VIRTUAL_NON_WORKER_CORES + 1 + 16;
constexpr static std::uint32_t MAX_NON_WORKER_CORES = 24 + 1 + 14;
constexpr static std::uint32_t MAX_HARVESTED_ROWS = 2;
constexpr static std::uint8_t CORE_COORD_INVALID = 0xFF;
struct core_info_msg_t {
Expand Down
4 changes: 4 additions & 0 deletions tt_metal/api/tt-metalium/device_impl.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -85,8 +85,12 @@ class Device : public IDevice {
// Ethernet API
CoreCoord ethernet_core_from_logical_core(const CoreCoord &logical_core) const override;
CoreCoord logical_core_from_ethernet_core(const CoreCoord &ethernet_core) const override;
// `skip_reserved_tunnel_cores` is ignored on BH because there are no ethernet cores used for Fast Dispatch
// tunneling
std::unordered_set<CoreCoord> get_active_ethernet_cores(bool skip_reserved_tunnel_cores=false) const override;
std::unordered_set<CoreCoord> get_inactive_ethernet_cores() const override;
// `skip_reserved_tunnel_cores` is ignored on BH because there are no ethernet cores used for Fast Dispatch
// tunneling
bool is_active_ethernet_core(CoreCoord logical_core, bool skip_reserved_tunnel_cores=false) const override;
std::tuple<chip_id_t, CoreCoord> get_connected_ethernet_core(CoreCoord eth_core) const override;
std::vector<CoreCoord> get_ethernet_sockets(chip_id_t connected_chip_id) const override;
Expand Down
2 changes: 2 additions & 0 deletions tt_metal/api/tt-metalium/hal.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -159,6 +159,7 @@ class Hal {
bool coordinate_virtualization_enabled_;
uint32_t virtual_worker_start_x_;
uint32_t virtual_worker_start_y_;
bool eth_fw_is_cooperative_; // set when eth riscs have to context switch

float eps_ = 0.0f;
float nan_ = 0.0f;
Expand Down Expand Up @@ -226,6 +227,7 @@ class Hal {
bool is_coordinate_virtualization_enabled() const { return this->coordinate_virtualization_enabled_; };
std::uint32_t get_virtual_worker_start_x() const { return this->virtual_worker_start_x_; }
std::uint32_t get_virtual_worker_start_y() const { return this->virtual_worker_start_y_; }
bool get_eth_fw_is_cooperative() const { return this->eth_fw_is_cooperative_; }
uint32_t get_programmable_core_type_count() const;
HalProgrammableCoreType get_programmable_core_type(uint32_t core_type_index) const;
uint32_t get_programmable_core_type_index(HalProgrammableCoreType programmable_core_type_index) const;
Expand Down
4 changes: 2 additions & 2 deletions tt_metal/common/metal_soc_descriptor.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -160,8 +160,8 @@ CoordSystem metal_SocDescriptor::get_umd_coord_system() const {
}

void metal_SocDescriptor::generate_logical_eth_coords_mapping() {
for (int i = 0; i < this->get_cores(CoreType::ETH).size(); i++) {
this->logical_eth_core_to_chan_map.insert({{0, i}, i});
for (const auto& logical_coord : this->get_cores(CoreType::ETH, CoordSystem::LOGICAL)) {
this->logical_eth_core_to_chan_map.insert({{logical_coord.x, logical_coord.y}, logical_coord.y});
}
}

Expand Down
2 changes: 2 additions & 0 deletions tt_metal/hw/firmware/src/active_erisc.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -121,6 +121,8 @@ int main() {

flush_erisc_icache();

firmware_config_init(mailboxes, ProgrammableCoreType::ACTIVE_ETH, DISPATCH_CLASS_ETH_DM0);

enum dispatch_core_processor_masks enables =
(enum dispatch_core_processor_masks)launch_msg_address->kernel_config.enables;

Expand Down
Loading

0 comments on commit 7409fd8

Please sign in to comment.