1D Fabric launch via device init #18712

aagarwalTT · 2025-03-06T15:23:01Z

Ticket

Link to Github Issue

Problem description

1D Fabric (EDM Fabric) is currently launched as a subdevice. This change moves the initialization of 1D Fabric to device init.

What's changed

Moved 1D Fabric implementation and relevant headers/files to metal from ttnn
Added the logic for launching 1D fabric during device init and corresponding sync and teardown sequence
Moved the 1D fabric namespace to the existing tt::tt_fabric namespace and got rid of tt::fabric

Things to be covered in follow-up PR(s):

Migrating the tests to use the new init flow and not subdevice
Renaming from EDM to 1D fabric wherever applicable

Checklist

All post commit CI passes
Blackhole Post commit CI passes (if applicable)
Model regression CI passes (if applicable)
Device performance regression CI passes (if applicable)
(For models and ops writers) Full new models tests CI passes (if applicable)
New/Existing tests provide coverage for changes

SeanNijjar · 2025-03-06T17:37:04Z

tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_1d_rx.cpp

General question: do we need dedicated/unique sender and receiver test kernels between 1D and 2D fabric if they aren't directly interacting with fabric? Can the 1D just be a simplified version (e.g. some compile time args to disable mcasting in 4 directions?

I think we can have common kernels for testing once we have some abstraction layer that unifies the two? currently the setup is bit different, so I created a new file to avoid confusion

SeanNijjar · 2025-03-06T17:38:56Z

tt_metal/api/tt-metalium/erisc_datamover_builder.hpp

This should be a git mv

It was initially, but then I split it to move the LineOpInterface into a helper file in ttnn itself

SeanNijjar

some initial comments. I'm not familiar with topology so I'll spend a little time later looking at it but would obviously defer to the other experts on the team for that design.

SeanNijjar · 2025-03-06T17:40:37Z

tt_metal/fabric/erisc_datamover_builder.cpp

SeanNijjar · 2025-03-06T17:41:43Z

tt_metal/fabric/hw/inc/edm_fabric/edm_fabric_worker_adapters.hpp

Is it possible to omit all the reformatting from this PR? It'll simplify review and keep it more incremental

I didnt update the clang format file, so I think this is the result of that. I will see if this can be reverted.

SeanNijjar · 2025-03-06T17:43:11Z

tt_metal/fabric/hw/inc/edm_fabric/fabric_erisc_datamover_channels.hpp

-    FORCE_INLINE bool all_eth_packets_completed() const {
-        return this->local_rdptr.is_caught_up_to(this->local_wrptr);
+    [[nodiscard]] FORCE_INLINE bool connection_is_live() const {
+        return *connection_live_semaphore == tt::fabric::EdmToEdmSender<0>::open_connection_value;


was this grabbed from a branch of mine or something? Unsure why it's needed in the PR but I think I did make this change elsewhere.

I am not sure how this change made it here, will have to dig a bit. I certainly did not pull anything from your branch though

I just realized that the order of functions was changed probably because of clang formatting

SeanNijjar · 2025-03-06T17:45:14Z

ttnn/cpp/ttnn/operations/ccl/erisc_datamover_builder_helper.cpp

Can you share the reason we need this here? What's the purpose?

Keeping this in ttnn would keep the current tests intact. We do not want this in metal since the fabric init has moved to device init. So until we port the tests, this can exist in ttnn as a helper to launch 1d fabric using sub device

omilyutin-tt · 2025-03-07T02:14:08Z

ttnn/cpp/ttnn/operations/ccl/erisc_datamover_builder_helper.hpp

+
+    // Device ID -> EDM Builders
+    std::unordered_map<size_t, std::vector<tt::tt_fabric::FabricEriscDatamoverBuilder>> edm_builders_forward_direction;
+    std::unordered_map<size_t, std::vector<tt::tt_fabric::FabricEriscDatamoverBuilder>> edm_builders_backward_direction;


Possible to make private?

There is a good chance that this file will be deprecated soon once the tests are migrated to the new init flow

omilyutin-tt · 2025-03-07T02:14:21Z

ttnn/cpp/ttnn/operations/ccl/erisc_datamover_builder_helper.hpp

+    // Device ID -> link index
+    std::unordered_map<size_t, size_t> next_forward_direction_edm_available;
+    std::unordered_map<size_t, size_t> next_backward_direction_edm_available;
+
+    std::vector<tt::tt_metal::IDevice*> device_sequence;
+    std::vector<tt::tt_metal::Program*> programs;
+
+    size_t num_links;
+    size_t buffer_size_bytes;
+    size_t firmware_context_switch_interval =
+        tt::tt_fabric::FabricEriscDatamoverBuilder::default_firmware_context_switch_interval;


Add suffix _ for data members

omilyutin-tt · 2025-03-07T02:14:31Z

ttnn/cpp/ttnn/operations/ccl/erisc_datamover_builder_helper.hpp

+    size_t num_links;
+    size_t buffer_size_bytes;


initialize to 0 to avoid UB

omilyutin-tt · 2025-03-07T02:14:54Z

ttnn/cpp/ttnn/operations/ccl/erisc_datamover_builder_helper.hpp

+        for (size_t i = 0; i < device_sequence.size(); i++) {
+            if (device_sequence[i] == device) {
+                return i;
+            }
+        }
+        TT_THROW("Device {} not found in device sequence of line fabric", device->id());


Possible to move to.cpp?

…e tests to capture missing updates

aliuTT · 2025-03-06T22:14:38Z

tt_metal/fabric/control_plane.cpp

+        if (routing_direction == direction) {
+            return eth_chans;
+        }
+    }


Would it be safer to assert here instead? Or are callers of get_active_fabric_eth_channels_in_direction responsible for checking for no eth channels in direction

aliuTT · 2025-03-08T03:58:00Z

tests/tt_metal/tt_fabric/common/fabric_fixture.hpp

@@ -83,4 +85,12 @@ class FabricFixture : public ::testing::Test {
    }
 };

+class FabricFixture : public BaseFabricFixture {


nit: can we call this Fabric2DFixture?

aagarwalTT changed the title ~~Aagarwal/1d fabric device init~~ 1d fabric device init Mar 6, 2025

SeanNijjar reviewed Mar 6, 2025

View reviewed changes

aagarwalTT changed the title ~~1d fabric device init~~ 1D Fabric launch via device init Mar 7, 2025

aagarwalTT marked this pull request as ready for review March 7, 2025 01:15

aagarwalTT requested review from ayerofieiev-tt, dmakoviichuk-tt, sminakov-tt, jvegaTT, cfjchu, tt-aho, ubcheema, aliuTT, abhullar-tt, pgkeller, tt-dma, tt-asaigal, davorchap and omilyutin-tt as code owners March 7, 2025 01:15

aagarwalTT force-pushed the aagarwal/1d-fabric-device-init branch from 87e09c3 to 265a94d Compare March 7, 2025 01:52

omilyutin-tt reviewed Mar 7, 2025

View reviewed changes

aagarwalTT added 5 commits March 7, 2025 08:51

moved files around, compiling and some tests passing, need to run mor…

fe9f6d4

…e tests to capture missing updates

missing files

f0f2844

device pool changes to launch 1d fabric

b1686df

add tests and some fixes

ffcc5ed

test updates and other fixes

c54a3ad

aagarwalTT added 4 commits March 7, 2025 08:51

namespace updates

142c4ad

minor updates

3593cc6

update clang-format-ignore to remove moved files

3be4a2c

ci fixes

a97630c

aagarwalTT force-pushed the aagarwal/1d-fabric-device-init branch from 95af811 to a97630c Compare March 7, 2025 14:51

aagarwalTT added 2 commits March 7, 2025 10:02

more ci fixes

17be215

cleanup and CI fixes

6976b04

ubcheema approved these changes Mar 7, 2025

View reviewed changes

aliuTT approved these changes Mar 8, 2025

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

1D Fabric launch via device init #18712

1D Fabric launch via device init #18712

aagarwalTT commented Mar 6, 2025 •

edited

Loading

SeanNijjar Mar 6, 2025

aagarwalTT Mar 6, 2025

SeanNijjar Mar 6, 2025

aagarwalTT Mar 6, 2025

SeanNijjar left a comment

SeanNijjar Mar 6, 2025

SeanNijjar Mar 6, 2025

aagarwalTT Mar 6, 2025

SeanNijjar Mar 6, 2025

aagarwalTT Mar 6, 2025

aagarwalTT Mar 7, 2025

SeanNijjar Mar 6, 2025

aagarwalTT Mar 6, 2025

omilyutin-tt Mar 7, 2025

aagarwalTT Mar 7, 2025 •

edited

Loading

omilyutin-tt Mar 7, 2025

omilyutin-tt Mar 7, 2025

omilyutin-tt Mar 7, 2025

aliuTT Mar 6, 2025

aliuTT Mar 8, 2025

1D Fabric launch via device init #18712

Are you sure you want to change the base?

1D Fabric launch via device init #18712

Conversation

aagarwalTT commented Mar 6, 2025 • edited Loading

Ticket

Problem description

What's changed

Checklist

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

SeanNijjar left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

aagarwalTT Mar 7, 2025 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

aagarwalTT commented Mar 6, 2025 •

edited

Loading

aagarwalTT Mar 7, 2025 •

edited

Loading