[HIP] fix(custom_all_reduce): uniform trip count for the block barriers

[Arist12]

What is wrong

Three kernels in csrc/include/custom_all_reduce.cuh stage their operands through shared memory
and bracket that with two block-wide __syncthreads() — inside a grid-stride loop whose trip
count is not uniform across the block:

kernel	loop
cross_device_reduce_2stage	stage 1, reduce-scatter
cross_device_reduce_2stage_write_mode	stage 2
reduce_scatter_cross_device_store	fused allreduce+rmsnorm, first step

Taking the first as the example:

for(int idx = start + tid; idx < end; idx += stride)
{
    *(reinterpret_cast<P*>(&tmp_smem[0]) + threadIdx.x) = ptrs[warp_id][idx];
    __syncthreads();
    if(warp_id == 0) { /* sum the ngpus operands out of smem */ }
    __syncthreads();
}

end - start is this rank's partition, size / ngpus, which is not generally a multiple of
stride = gridDim.x * (THREAD_NUM / ngpus). The lanes that cover the tail therefore run one
iteration more than the rest, so the barriers pair up across different iterations — and against
the ones in end_sync. Two consequences:

• the warp_id == 0 reduction can read tmp_smem before the other wavefronts have written
  their operands into it;
• a wavefront can reach end_sync and publish "stage 1 done" while another is still producing
  tmp_out, so the stage-2 all-gather can read a partition that is not finished.

Either way the output is wrong at a fixed relative position inside each rank's partition.

How it shows up

size is flat, so each rank owns a contiguous size / ngpus slice — one quarter of the rows of
a [bs, hidden] tensor at ngpus == 4. The damage lands at the same offset in every quarter:

bs = 120  ->  bad rows {1, 31, 61, 91}      (stride 30 = 120/4)
bs = 160  ->  bad rows {11, 51, 91, 131}    (stride 40 = 160/4)

Found in an RL workload; reproduced standalone below.

Standalone reproduction

Four ranks on MI355X, bf16 [bs, 2880], 32 shapes × 600 reductions each. Every reduction is run
twice and compared bitwise (no tolerance needed), and also compared against RCCL's all_reduce
(bf16 rounding on a 4-way sum is ~0.05, so anything above that is a wrong answer rather than a
different summation order):

	before	after
wrong vs RCCL, and non-deterministic	64 of 19200	0 of 19200
shapes affected	bs = 120, 121, 127, 157, 160, 161	none
worst |diff|	12.4	0.0625 (pure bf16 rounding)
48 reductions back to back, bs=160	300 of 11520 wrong, worst 13.1	0 of 11520
world_size=2 (1-stage path)	0 of 5120	0 of 5120 — no regression

It is a race, so the counts move between runs; an earlier sweep over a slightly different set of
shapes gave 43/19200 and 173/11520. Most shapes pass either way — the race needs the partition
remainder to land mid-block, which is why this has survived.

The numbers above were measured on a build based on d9e5ef7ce that also carried a second
commit, since dropped, which widened two sync scopes. @TennyWang1223 re-measured 6110f00f
on its own and reproduced the fix, which attributes it to the barrier change alone. The
three patched hunks are unchanged on main, and the rebased diff is byte-identical to the
one that was tested.

The fix

Round the loop bound up to a multiple of the stride, which makes the trip count exactly uniform
(tid < stride by construction), and predicate the work. No added synchronisation and no change
to the reduction order.

reduce_scatter_cross_device_store is fixed by inspection — same defect, same file — but the
reproducer above does not exercise the fused-allreduce path.

Timing, 7 reps × 500 calls, median us/call taken as the max over ranks, whole measurement
repeated twice: within run-to-run noise at bs=120..161, and 3-6% slower at bs=256
(19.2 → 20.3) and bs=512 (32.2 → 33.1).

[github-actions]

🏷️ CI Guide

Runs automatically on every PR:

• ✅ Pre-checks (submodule verification, code formatting)
• ✅ Aiter op tests (gfx942 + gfx950)
• ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label	Tests
ci:gfx1250-ffm-triton	Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x	Run an additional Triton test job on MI300X in PRs; main branch always runs both MI35X and MI300X
multigpu	Aiter multi-GPU tests on the 8-GPU runner
ci:sglang	SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom	ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full	ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm	vLLM benchmark: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
ci:all	All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 5694 --add-label <label>

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to opt this PR out.

[TennyWang1223]

1. The change of device function start_sync and end_sync are unnecessary. detail see #4294 (comment)
2. The change on kernels are truely useful, I have tested on main and this pr without change start_sync and end_sync(commit 6110f00), results as bellow. this can also prove origin start_sync and end_sync have no influence on accuracy

main

commit 6110f00

[Arist12]

1. The change of device function start_sync and end_sync are unnecessary. detail see #4294 (comment)
2. The change on kernels are truely useful, I have tested on main and this pr without change start_sync and end_sync(commit 6110f00), results as bellow. this can also prove origin start_sync and end_sync have no influence on accuracy

main

commit 6110f00

1. The change of device function start_sync and end_sync are unnecessary. detail see #4294 (comment)
2. The change on kernels are truely useful, I have tested on main and this pr without change start_sync and end_sync(commit 6110f00), results as bellow. this can also prove origin start_sync and end_sync have no influence on accuracy

main

commit 6110f00

hi @TennyWang1223 , thanks for the careful review, I've trimmed the corresponding section per requested.

[TennyWang1223]

LGTM

[Arist12]

The new commit is syntax-only, only trim the comments and the code should function well; should be good to merge @TennyWang1223