Skip to content

Commit 84de900

Browse files
tt-asaigaltt-asaigal
authored
Update EDM <-> EDM handshake in fabric_erisc_datamover (#23813)
### Ticket No Ticket. ### Problem description - The current EDM <-> EDM handshake approach for `fabric_erisc_datamover` kernels relies on all devices in the Mesh being initialized before fabric is loaded - This assumption is true for single host scenarios, but breaks down in the multi-host case, since each host will initialize it's local mesh asynchronous of all other hosts - With delays added to device startup on some hosts, the handshake step hangs. This is because the delayed host clears the handshake message written to it ### What's changed - Add new _single step_ handshake APIs to `edm_handshake.hpp`: `sender_side_handshake` and `receiver_side_handshake` (protocol described in header) - Use them in `fabric_erisc_datamover.cpp` to ensure that the multi-host scenario described above is supported - Mark handshake APIs used by legacy CCLs as deprecated. These will be removed once legacy CCLs are purged. - General cleanup - remove unused args from `EthChannelBuffer` ctor ### Checklist - [ ] [All post commit](https://github.com/tenstorrent/tt-metal/actions/workflows/all-post-commit-workflows.yaml) CI passes - [ ] [Blackhole Post commit](https://github.com/tenstorrent/tt-metal/actions/workflows/blackhole-post-commit.yaml) CI with demo tests passes (if applicable) - [ ] [Model regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-models.yaml) CI passes (if applicable) - [ ] [Device performance regression](https://github.com/tenstorrent/tt-metal/actions/workflows/perf-device-models.yaml) CI passes (if applicable) - [ ] (For models and ops writers) [Single-card demo tests](https://github.com/tenstorrent/tt-metal/actions/workflows/single-card-demo-tests.yaml) CI passes (if applicable) See [recommended dev flow](https://github.com/tenstorrent/tt-metal/blob/main/models/MODEL_ADD.md#a-recommended-dev-flow-on-github-for-adding-new-models). - [ ] New/Existing tests provide coverage for changes
1 parent 30110bc commit 84de900

File tree

10 files changed

+111
-60
lines changed

10 files changed

+111
-60
lines changed

‎tests/tt_metal/tt_metal/perf_microbenchmark/routing/kernels/tt_fabric_mux_ubench_drainer.cpp‎

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -40,11 +40,7 @@ void kernel_main() {
4040
init_ptr_val(slots_free_stream_id, NUM_BUFFERS);
4141

4242
tt::tt_fabric::DrainerChannelBuffer drainer_channel(
43-
channel_base_address,
44-
BUFFER_SIZE_BYTES,
45-
sizeof(PACKET_HEADER_TYPE),
46-
0, /* unused, eth_transaction_ack_word_addr */
47-
0 /* channel_id */);
43+
channel_base_address, BUFFER_SIZE_BYTES, sizeof(PACKET_HEADER_TYPE), 0 /* channel_id */);
4844

4945
auto connection_worker_info_ptr =
5046
reinterpret_cast<volatile tt::tt_fabric::DrainerChannelClientLocationInfo*>(connection_info_address);

‎tt_metal/fabric/erisc_datamover_builder.cpp‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -808,7 +808,10 @@ std::vector<uint32_t> FabricEriscDatamoverBuilder::get_compile_time_args(uint32_
808808

809809
// TODO: allow specification per eth txq
810810
const size_t default_num_eth_txq_data_packet_accept_ahead = 32;
811-
811+
// By default have the ERISC cores context switch to base routing FW every 4K cycles during the peer handshake.
812+
// This allows host to write Fabric kernels to remote chips over ethernet, when ERISC cores already running fabric
813+
// are waiting for the handshake to complete.
814+
const size_t default_handshake_context_switch_timeout = 4096;
812815
size_t num_sender_channels = config.num_used_sender_channels;
813816
size_t num_receiver_channels = config.num_used_receiver_channels;
814817
const auto& control_plane = tt::tt_metal::MetalContext::instance().get_control_plane();
@@ -915,6 +918,7 @@ std::vector<uint32_t> FabricEriscDatamoverBuilder::get_compile_time_args(uint32_
915918
eth_txq_spin_wait_receiver_send_completion_ack,
916919
default_num_eth_txq_data_packet_accept_ahead,
917920

921+
default_handshake_context_switch_timeout,
918922
// Special marker to help with identifying misalignment bugs
919923
0x00c0ffee};
920924

‎tt_metal/fabric/hw/inc/edm_fabric/1d_fabric_constants.hpp‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -295,7 +295,10 @@ constexpr bool ETH_TXQ_SPIN_WAIT_RECEIVER_SEND_COMPLETION_ACK = get_compile_time
295295

296296
constexpr size_t DEFAULT_NUM_ETH_TXQ_DATA_PACKET_ACCEPT_AHEAD = get_compile_time_arg_val(MAIN_CT_ARGS_IDX_5 + 11);
297297

298-
constexpr size_t SPECIAL_MARKER_0_IDX = MAIN_CT_ARGS_IDX_5 + 12;
298+
// Context switch timeouts
299+
constexpr size_t DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT = get_compile_time_arg_val(MAIN_CT_ARGS_IDX_5 + 12);
300+
301+
constexpr size_t SPECIAL_MARKER_0_IDX = MAIN_CT_ARGS_IDX_5 + 13;
299302
constexpr size_t SPECIAL_MARKER_0 = 0x00c0ffee;
300303
static_assert(
301304
!SPECIAL_MARKER_CHECK_ENABLED || get_compile_time_arg_val(SPECIAL_MARKER_0_IDX) == SPECIAL_MARKER_0,
@@ -382,7 +385,6 @@ constexpr std::array<uint32_t, MAX_NUM_SENDER_CHANNELS> to_sender_packets_comple
382385
to_sender_3_pkts_completed_id, to_sender_4_pkts_completed_id});
383386

384387
// Miscellaneous configuration
385-
constexpr size_t DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT = 0;
386388

387389
// TODO: move this to compile time args if we need to enable it
388390
constexpr bool enable_trid_flush_check_on_noc_txn = false;

‎tt_metal/fabric/hw/inc/edm_fabric/edm_handshake.hpp‎

Lines changed: 84 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66

77
#include <cstdint>
88
#include "ethernet/dataflow_api.h"
9+
#include <fabric_host_interface.h>
910

1011
namespace erisc {
1112
namespace datamover {
@@ -16,11 +17,6 @@ namespace datamover {
1617
* case. Before handshaking, we make sure to clear any of the channel sync datastructures local
1718
* to our core.
1819
*
19-
* The handshaking process is split into two parts for the sender/master and two parts for the
20-
* the subordinate. The handshake is broken into 2 parts so that the master can initiate the handshake
21-
* as early as possible so the message can be "in flight" over the ethernet link while other EDM
22-
* initialization is taking place.
23-
*
2420
* Important note about handshaking: the sender/master canNOT complete the handshake until all receiver
2521
* channels are initialized. Otherwise we have a race between channel initialization on the receiver side
2622
* and real payload data (and signals) using those channels.
@@ -32,9 +28,84 @@ namespace datamover {
3228
* and sends payload available information to that channel. The receive must acknowledge that message and upon
3329
* doing so, considers the handshake complete.
3430
*/
31+
3532
namespace handshake {
3633

34+
/* EDM Handshaking Mechanism:
35+
* 1. Both sides set their local_value register to 0.
36+
* 2. Both sides write a magic value to their scratch register.
37+
* 3. Handshake master repeatedly copies the magic value from the scratch register to the local_value of the remote
38+
* subordinate, until it sees the magic value in its local_value register.
39+
* 4. Handshake subordinate polls its local_value register until it sees the magic value written by the master. It
40+
* then copies the magic value from its scratch register to the master's local_value register, completing the handshake.
41+
*/
42+
3743
static constexpr uint32_t A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH = 1000000000;
44+
static constexpr uint32_t MAGIC_HANDSHAKE_VALUE = 0xAA;
45+
46+
// Data-Structure used for EDM to EDM Handshaking.
47+
struct handshake_info_t {
48+
uint32_t local_value; // Updated by remote
49+
uint32_t padding[3]; // Ensures 16B alignment for scratch register
50+
uint32_t scratch[4]; // TODO: This can be removed if we use a stream register for handshaking.
51+
};
52+
53+
FORCE_INLINE volatile tt_l1_ptr handshake_info_t* init_handshake_info(uint32_t handshake_register_address) {
54+
volatile tt_l1_ptr handshake_info_t* handshake_info =
55+
reinterpret_cast<volatile tt_l1_ptr handshake_info_t*>(handshake_register_address);
56+
handshake_info->local_value = 0;
57+
handshake_info->scratch[0] = MAGIC_HANDSHAKE_VALUE;
58+
return handshake_info;
59+
}
60+
61+
FORCE_INLINE void sender_side_handshake(
62+
uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
63+
volatile tt_l1_ptr handshake_info_t* handshake_info = init_handshake_info(handshake_register_address);
64+
uint32_t local_val_addr = ((uint32_t)(&handshake_info->local_value)) / tt::tt_fabric::PACKET_WORD_SIZE_BYTES;
65+
uint32_t scratch_addr = ((uint32_t)(&handshake_info->scratch)) / tt::tt_fabric::PACKET_WORD_SIZE_BYTES;
66+
uint32_t count = 0;
67+
while (handshake_info->local_value != MAGIC_HANDSHAKE_VALUE) {
68+
if (count == HS_CONTEXT_SWITCH_TIMEOUT) {
69+
count = 0;
70+
run_routing();
71+
} else {
72+
count++;
73+
internal_::eth_send_packet(0, scratch_addr, local_val_addr, 1);
74+
}
75+
invalidate_l1_cache();
76+
}
77+
}
78+
79+
FORCE_INLINE void receiver_side_handshake(
80+
uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
81+
volatile tt_l1_ptr handshake_info_t* handshake_info = init_handshake_info(handshake_register_address);
82+
uint32_t local_val_addr = ((uint32_t)(&handshake_info->local_value)) / tt::tt_fabric::PACKET_WORD_SIZE_BYTES;
83+
uint32_t scratch_addr = ((uint32_t)(&handshake_info->scratch)) / tt::tt_fabric::PACKET_WORD_SIZE_BYTES;
84+
uint32_t count = 0;
85+
while (handshake_info->local_value != MAGIC_HANDSHAKE_VALUE) {
86+
if (count == HS_CONTEXT_SWITCH_TIMEOUT) {
87+
count = 0;
88+
run_routing();
89+
} else {
90+
count++;
91+
}
92+
invalidate_l1_cache();
93+
}
94+
internal_::eth_send_packet(0, scratch_addr, local_val_addr, 1);
95+
}
96+
97+
namespace deprecated {
98+
99+
/* The split handshaking mechanism exposed through the following APIs is deprecated.
100+
* This was developed for non-persistent kernels targeting EDM cores, which are not
101+
* supported with TT-Fabric.
102+
* TODO: Remove these APIs once legacy CCL Ops are removed.
103+
*
104+
* The handshaking process is split into two parts for the sender/master and two parts for the
105+
* the subordinate. The handshake is broken into 2 parts so that the master can initiate the handshake
106+
* as early as possible so the message can be "in flight" over the ethernet link while other EDM
107+
* initialization is taking place.
108+
*/
38109

39110
/*
40111
* Initialize base datastructures and values which are common to master and subordinate EDM cores.
@@ -47,7 +118,7 @@ static constexpr uint32_t A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH = 1000000000;
47118
*
48119
* See ChannelBuffer::eth_receiver_channel_ack for more information
49120
*/
50-
FORCE_INLINE void initialize_edm_common_datastructures(std::uint32_t handshake_register_address) {
121+
FORCE_INLINE void initialize_edm_common_datastructures(uint32_t handshake_register_address) {
51122
reinterpret_cast<volatile tt_l1_ptr uint32_t*>(handshake_register_address)[4] = 1;
52123
reinterpret_cast<volatile tt_l1_ptr uint32_t*>(handshake_register_address)[5] = 1;
53124
reinterpret_cast<volatile tt_l1_ptr uint32_t*>(handshake_register_address)[6] = 0x1c0ffee1;
@@ -66,7 +137,7 @@ FORCE_INLINE void initialize_edm_common_datastructures(std::uint32_t handshake_r
66137
* memory region.
67138
*/
68139
FORCE_INLINE void sender_side_start(
69-
std::uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
140+
uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
70141
initialize_edm_common_datastructures(handshake_register_address);
71142
eth_wait_receiver_done(HS_CONTEXT_SWITCH_TIMEOUT);
72143
while (eth_txq_is_busy()) {
@@ -79,11 +150,11 @@ FORCE_INLINE void sender_side_start(
79150
* As the designated master EDM core, wait for the acknowledgement from the subordinate EDM core
80151
*/
81152
FORCE_INLINE void sender_side_finish(
82-
std::uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
153+
uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
83154
eth_wait_for_receiver_done(HS_CONTEXT_SWITCH_TIMEOUT);
84155
}
85156

86-
FORCE_INLINE void receiver_side_start(std::uint32_t handshake_register_address) {
157+
FORCE_INLINE void receiver_side_start(uint32_t handshake_register_address) {
87158
initialize_edm_common_datastructures(handshake_register_address);
88159
}
89160

@@ -99,13 +170,16 @@ FORCE_INLINE bool receiver_side_can_finish() { return eth_bytes_are_available_on
99170
* from the master EDM core.
100171
*/
101172
FORCE_INLINE void receiver_side_finish(
102-
std::uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
173+
uint32_t handshake_register_address, size_t HS_CONTEXT_SWITCH_TIMEOUT = A_LONG_TIMEOUT_BEFORE_CONTEXT_SWITCH) {
103174
eth_wait_for_bytes(16, HS_CONTEXT_SWITCH_TIMEOUT);
104175
while (eth_txq_is_busy()) {
105176
asm volatile("nop");
106177
}
107178
eth_receiver_channel_done(0);
108179
}
180+
181+
} // namespace deprecated
182+
109183
} // namespace handshake
110184

111185
} // namespace datamover

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

Lines changed: 1 addition & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -53,12 +53,7 @@ class EthChannelBuffer final {
5353
* Expected that *buffer_index_ptr is initialized outside of this object
5454
*/
5555
EthChannelBuffer(
56-
size_t channel_base_address,
57-
size_t buffer_size_bytes,
58-
size_t header_size_bytes,
59-
size_t eth_transaction_ack_word_addr, // Assume for receiver channel, this address points to a chunk of memory
60-
// that can fit 2 eth_channel_syncs cfor ack
61-
uint8_t channel_id) :
56+
size_t channel_base_address, size_t buffer_size_bytes, size_t header_size_bytes, uint8_t channel_id) :
6257
buffer_size_in_bytes(buffer_size_bytes),
6358
max_eth_payload_size_in_bytes(buffer_size_in_bytes),
6459
channel_id(channel_id) {
@@ -134,7 +129,6 @@ struct EthChannelBufferTuple {
134129
const size_t channel_base_address[],
135130
const size_t buffer_size_bytes,
136131
const size_t header_size_bytes,
137-
const size_t eth_transaction_ack_word_addr,
138132
const size_t channel_base_id) {
139133
size_t idx = 0;
140134

@@ -144,7 +138,6 @@ struct EthChannelBufferTuple {
144138
channel_base_address[idx],
145139
buffer_size_bytes,
146140
header_size_bytes,
147-
eth_transaction_ack_word_addr,
148141
static_cast<uint8_t>(channel_base_id + idx)),
149142
++idx),
150143
...);

‎tt_metal/fabric/impl/kernels/edm_fabric/fabric_erisc_datamover.cpp‎

Lines changed: 5 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -1425,8 +1425,6 @@ void kernel_main() {
14251425
//
14261426
// COMMON CT ARGS (not specific to sender or receiver)
14271427
//
1428-
*reinterpret_cast<volatile uint32_t*>(handshake_addr) = 0;
1429-
auto eth_transaction_ack_word_addr = handshake_addr + sizeof(eth_channel_sync_t);
14301428

14311429
// Initialize stream register state for credit management across the Ethernet link.
14321430
// We make sure to do this before we handshake to guarantee that the registers are
@@ -1466,14 +1464,6 @@ void kernel_main() {
14661464
init_ptr_val<to_sender_packets_completed_streams[4]>(0);
14671465
}
14681466

1469-
if constexpr (enable_ethernet_handshake) {
1470-
if constexpr (is_handshake_sender) {
1471-
erisc::datamover::handshake::sender_side_start(handshake_addr, DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT);
1472-
} else {
1473-
erisc::datamover::handshake::receiver_side_start(handshake_addr);
1474-
}
1475-
}
1476-
14771467
// TODO: CONVERT TO SEMAPHORE
14781468
volatile auto termination_signal_ptr =
14791469
reinterpret_cast<volatile tt::tt_fabric::TerminationSignal*>(termination_signal_addr);
@@ -1875,24 +1865,18 @@ void kernel_main() {
18751865
local_receiver_buffer_addresses.data(),
18761866
channel_buffer_size,
18771867
sizeof(PACKET_HEADER_TYPE),
1878-
eth_transaction_ack_word_addr,
18791868
receiver_channel_base_id);
18801869

18811870
// initialize the remote receiver channel buffers
18821871
remote_receiver_channels.init(
18831872
remote_receiver_buffer_addresses.data(),
18841873
channel_buffer_size,
18851874
sizeof(PACKET_HEADER_TYPE),
1886-
eth_transaction_ack_word_addr,
18871875
receiver_channel_base_id);
18881876

18891877
// initialize the local sender channel worker interfaces
18901878
local_sender_channels.init(
1891-
local_sender_buffer_addresses.data(),
1892-
channel_buffer_size,
1893-
sizeof(PACKET_HEADER_TYPE),
1894-
0, // For sender channels there is no eth_transaction_ack_word_addr because they don't send acks
1895-
sender_channel_base_id);
1879+
local_sender_buffer_addresses.data(), channel_buffer_size, sizeof(PACKET_HEADER_TYPE), sender_channel_base_id);
18961880

18971881
// initialize the local sender channel worker interfaces
18981882
init_local_sender_channel_worker_interfaces(
@@ -1921,9 +1905,11 @@ void kernel_main() {
19211905

19221906
if constexpr (enable_ethernet_handshake) {
19231907
if constexpr (is_handshake_sender) {
1924-
erisc::datamover::handshake::sender_side_finish(handshake_addr, DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT);
1908+
erisc::datamover::handshake::sender_side_handshake(
1909+
handshake_addr, DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT);
19251910
} else {
1926-
erisc::datamover::handshake::receiver_side_finish(handshake_addr, DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT);
1911+
erisc::datamover::handshake::receiver_side_handshake(
1912+
handshake_addr, DEFAULT_HANDSHAKE_CONTEXT_SWITCH_TIMEOUT);
19271913
}
19281914

19291915
*edm_status_ptr = tt::tt_fabric::EDMStatus::REMOTE_HANDSHAKE_COMPLETE;

‎tt_metal/fabric/impl/kernels/tt_fabric_mux.cpp‎

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -55,11 +55,7 @@ void setup_channel(
5555
size_t& sender_flow_control_address,
5656
StreamId my_channel_free_slots_stream_id) {
5757
new (channel_ptr) tt::tt_fabric::FabricMuxChannelBuffer<NUM_BUFFERS>(
58-
channel_base_address,
59-
buffer_size_bytes,
60-
sizeof(PACKET_HEADER_TYPE),
61-
0, /* unused, eth_transaction_ack_word_addr */
62-
channel_id);
58+
channel_base_address, buffer_size_bytes, sizeof(PACKET_HEADER_TYPE), channel_id);
6359
channel_base_address += NUM_BUFFERS * buffer_size_bytes;
6460
init_ptr_val(my_channel_free_slots_stream_id, NUM_BUFFERS);
6561

‎ttnn/cpp/ttnn/operations/ccl/barrier/device/kernels/barrier_receiver.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -55,8 +55,8 @@ void kernel_main() {
5555
channels_syncs_addrs->bytes_sent = 0;
5656
channels_syncs_addrs->receiver_ack = 0;
5757
// Semaphore is mapped to sender core
58-
erisc::datamover::handshake::receiver_side_start(handshake_addr);
59-
erisc::datamover::handshake::receiver_side_finish(handshake_addr);
58+
erisc::datamover::handshake::deprecated::receiver_side_start(handshake_addr);
59+
erisc::datamover::handshake::deprecated::receiver_side_finish(handshake_addr);
6060

6161
*start_semaphore = 0;
6262

‎ttnn/cpp/ttnn/operations/ccl/barrier/device/kernels/barrier_sender.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ void kernel_main() {
4646
uint32_t arg_idx = 0;
4747
const bool is_ring_start = get_arg_val<uint32_t>(arg_idx++) == 1;
4848
const uint32_t handshake_addr = get_arg_val<uint32_t>(arg_idx++);
49-
erisc::datamover::handshake::sender_side_start(handshake_addr);
49+
erisc::datamover::handshake::deprecated::sender_side_start(handshake_addr);
5050
uint32_t channels_addrs = get_arg_val<uint32_t>(arg_idx++);
5151
volatile uint32_t* sem_addr = reinterpret_cast<volatile uint32_t*>(get_arg_val<uint32_t>(arg_idx++));
5252
const uint32_t host_noc_x = get_arg_val<uint32_t>(arg_idx++);
@@ -60,7 +60,7 @@ void kernel_main() {
6060
channels_syncs_addrs->bytes_sent = 0;
6161
channels_syncs_addrs->receiver_ack = 0;
6262
*sem_addr = 0;
63-
erisc::datamover::handshake::sender_side_finish(handshake_addr);
63+
erisc::datamover::handshake::deprecated::sender_side_finish(handshake_addr);
6464
noc_semaphore_inc(host_semaphore_addr, 1);
6565

6666
// Ensure every core has completed their previous tasks

‎ttnn/cpp/ttnn/operations/ccl/kernels/edm/erisc_datamover.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -145,9 +145,9 @@ void kernel_main() {
145145

146146
bool is_done_as_rx_handshaker = is_handshake_sender;
147147
if constexpr (is_handshake_sender) {
148-
erisc::datamover::handshake::sender_side_start(handshake_addr);
148+
erisc::datamover::handshake::deprecated::sender_side_start(handshake_addr);
149149
} else {
150-
erisc::datamover::handshake::receiver_side_start(handshake_addr);
150+
erisc::datamover::handshake::deprecated::receiver_side_start(handshake_addr);
151151
}
152152

153153
// Receiver args
@@ -184,9 +184,9 @@ void kernel_main() {
184184
}
185185

186186
if (!is_handshake_sender) {
187-
if (!is_done_as_rx_handshaker && erisc::datamover::handshake::receiver_side_can_finish()) {
187+
if (!is_done_as_rx_handshaker && erisc::datamover::handshake::deprecated::receiver_side_can_finish()) {
188188
is_done_as_rx_handshaker = true;
189-
erisc::datamover::handshake::receiver_side_finish(handshake_addr);
189+
erisc::datamover::handshake::deprecated::receiver_side_finish(handshake_addr);
190190
}
191191
}
192192

@@ -224,10 +224,10 @@ void kernel_main() {
224224
}
225225

226226
if constexpr (is_handshake_sender) {
227-
erisc::datamover::handshake::sender_side_finish(handshake_addr);
227+
erisc::datamover::handshake::deprecated::sender_side_finish(handshake_addr);
228228
} else {
229229
if (!is_done_as_rx_handshaker) {
230-
erisc::datamover::handshake::receiver_side_finish(handshake_addr);
230+
erisc::datamover::handshake::deprecated::receiver_side_finish(handshake_addr);
231231
is_done_as_rx_handshaker = true;
232232
}
233233
}

0 commit comments

Comments
 (0)