[ROCm] Fix intranode queue ordering for shared buffers

[AndreasKaratzas]

The Distributed Tests / DBO test group on ROCm vLLM CI fails in build 11147 when successive DBO microbatches share one ROCm DeepEP buffer. The intranode path can publish queue progress without a complete producer/consumer ownership chain, while whole-workgroup barriers can pair unrelated per-rank rounds. This change uses per-rank wave barriers and release/acquire ordering for dispatch and combine tail publication, queue-slot reuse, receiver progress, retirement, and shared tail fan-out. It also adds a two-rank regression test that drives asymmetric and randomized microbatches through repeated queue wraparound on the same buffer.

[kudomcho]

Do you have test suite results after this change?

[AndreasKaratzas]

@kudomcho For vLLM you mean?

[AndreasKaratzas]

@kudomcho lmk if there is anything needed on my end :)

[kudomcho]

Hi, yes, please provide test reproducer and result logs if possible. Thank you.

[AndreasKaratzas]

So the failure is:

• https://buildkite.com/vllm/amd-ci/builds/11147/list?jid=019f8bae-0610-4768-8be3-65983033438f&tab=output#L8542

It is from our CI. You can use:

docker pull rocm/vllm-ci:da3a252fd13f51c22657bfc8650936f2fbb5b6f3

Which is one of our latest nightlies. And then you can just:

cd /vllm-workspace
pytest tests/v1/distributed/test_dbo.py::test_dbo_dp_ep_gsm8k[deepep_high_throughput]

You should then get a similar error message to:

2026-07-22 17:00:51 CDT
_________________________________________________________ test_dbo_dp_ep_gsm8k[deepep_high_throughput] _________________________________________________________
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
all2all_backend = 'deepep_high_throughput', num_gpus_available = 2
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
    @pytest.mark.skipif(not has_deep_ep(), reason="These tests require deep_ep to run")
2026-07-22 17:00:51 CDT
    @pytest.mark.parametrize("all2all_backend", DEEPEP_BACKENDS)
2026-07-22 17:00:51 CDT
    @pytest.mark.xfail(
2026-07-22 17:00:51 CDT
        IS_BLACKWELL,
2026-07-22 17:00:51 CDT
        reason=(
2026-07-22 17:00:51 CDT
            "Temporary: DBO accuracy unstable on Blackwell "
2026-07-22 17:00:51 CDT
            "(doesn't meet expectation of MIN_ACCURACY = 0.62)"
2026-07-22 17:00:51 CDT
        ),
2026-07-22 17:00:51 CDT
    )
2026-07-22 17:00:51 CDT
    def test_dbo_dp_ep_gsm8k(all2all_backend: str, num_gpus_available):
2026-07-22 17:00:51 CDT
        """
2026-07-22 17:00:51 CDT
        Test DBO with DP+EP using GSM8K evaluation.
2026-07-22 17:00:51 CDT
        """
2026-07-22 17:00:51 CDT
        required_gpus = DP_SIZE
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
        if num_gpus_available < required_gpus:
2026-07-22 17:00:51 CDT
            pytest.skip(f"Need at least {required_gpus} GPUs (DP={DP_SIZE})")
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
        # Server arguments for DBO + DP + EP
2026-07-22 17:00:51 CDT
        server_args = [
2026-07-22 17:00:51 CDT
            "--max-model-len",
2026-07-22 17:00:51 CDT
            "4096",
2026-07-22 17:00:51 CDT
            "--max-num-seqs",
2026-07-22 17:00:51 CDT
            str(MAX_NUM_SEQS),  # Use larger batch to trigger decode DBO
2026-07-22 17:00:51 CDT
            "--trust-remote-code",
2026-07-22 17:00:51 CDT
            # Note: Not using --enforce-eager to test DBO's alternate CUDA graph dispatching
2026-07-22 17:00:51 CDT
            "--data-parallel-size",
2026-07-22 17:00:51 CDT
            str(DP_SIZE),
2026-07-22 17:00:51 CDT
            "--enable-expert-parallel",
2026-07-22 17:00:51 CDT
            "--enable-dbo",
2026-07-22 17:00:51 CDT
            # Fix threshold so we know we trigger DBO
2026-07-22 17:00:51 CDT
            "--dbo-decode-token-threshold",
2026-07-22 17:00:51 CDT
            "16",
2026-07-22 17:00:51 CDT
            "--dbo-prefill-token-threshold",
2026-07-22 17:00:51 CDT
            "256",
2026-07-22 17:00:51 CDT
            "--all2all-backend",
2026-07-22 17:00:51 CDT
            all2all_backend,
2026-07-22 17:00:51 CDT
        ]
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
        with RemoteOpenAIServer(
2026-07-22 17:00:51 CDT
            MODEL_NAME,
2026-07-22 17:00:51 CDT
            server_args,
2026-07-22 17:00:51 CDT
            max_wait_seconds=600,  # Allow time for model loading with DP+EP
2026-07-22 17:00:51 CDT
        ) as remote_server:
2026-07-22 17:00:51 CDT
            # Use host and port directly from RemoteOpenAIServer
2026-07-22 17:00:51 CDT
            host = f"http://{remote_server.host}"
2026-07-22 17:00:51 CDT
            port = remote_server.port
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
            # Run GSM8K evaluation
2026-07-22 17:00:51 CDT
            results = evaluate_gsm8k(
2026-07-22 17:00:51 CDT
                num_questions=NUM_QUESTIONS,
2026-07-22 17:00:51 CDT
                num_shots=NUM_SHOTS,
2026-07-22 17:00:51 CDT
                host=host,
2026-07-22 17:00:51 CDT
                port=port,
2026-07-22 17:00:51 CDT
            )
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
            # Validate accuracy is reasonable
2026-07-22 17:00:51 CDT
            accuracy = results["accuracy"]
2026-07-22 17:00:51 CDT
>           assert accuracy >= MIN_ACCURACY, (
2026-07-22 17:00:51 CDT
                f"DBO+DP+EP accuracy too low ({all2all_backend}): "
2026-07-22 17:00:51 CDT
                f"{accuracy:.3f} < {MIN_ACCURACY:.3f} "
2026-07-22 17:00:51 CDT
            )
2026-07-22 17:00:51 CDT
E           AssertionError: DBO+DP+EP accuracy too low (deepep_high_throughput): 0.559 < 0.620 
2026-07-22 17:00:51 CDT
E           assert np.float64(0.55859375) >= 0.62
2026-07-22 17:00:51 CDT
2026-07-22 17:00:51 CDT
tests/v1/distributed/test_dbo.py:106: AssertionError

[AndreasKaratzas]

cc @kudomcho @avbokovoy @amirakb89

[amirakb89]

Thanks for contributing to the project, I was not able to reproduce the issue on MI300 with the pipeline_test.
I see the PR is closed, a few points that might help the future investigation:
1- similar wide-warp sync (bar.sync) to wait_workgroup_warp_barrier has been tested before and it found that they cause performance regression. They should be avoided as much as possible.
2- On MI300/XGMI the load/store semantics can be relaxed as long as the buffers are not being updated by different CUs.