Question: Why did DeepEP v2 drop the double-buffer design in favor of explicit pre/post-sync barriers?

Hi,

I'm comparing DeepEP v1 (legacy) and v2, and noticed a major change in the synchronization and buffering strategy.

In v1, the `low_latency` path avoided the `notify` pre-sync kernels by using a double-buffer (ping-pong) design combined with in-kernel atomic/flag polling over NVLink to overlap communication and computation.

However, in v2, it seems this double-buffer design is dropped, and the implementation goes back to using explicit pre-sync and post-sync barriers.

Could you share the design trade-offs behind this change? Specifically:

1. Did the double-buffer design in v1 LL introduce bottleneck issues in practice (e.g., memory/L2 cache pressure, or NVLink congestion from atomic polling)?
2. Does v2's explicit barrier leverage Hopper/Blackwell hardware features (like `mbarrier` or cluster-level sync) to make the barrier overhead negligible?
3. How does v2 handle low-latency scenarios without the ping-pong overlap?

Thanks!

Compare the DeepEP v1 and v2 low_latency implementations, focusing on the double-buffer and atomic-polling path versus the explicit pre-sync and post-sync barriers. Trace how each version handles overlap and identify the hardware synchronization features or measured trade-offs needed to answer the three questions.

Written by the indexing model from the issue text.