[hybrid_dispatch] Question on release scope of local-rank Rail bypass path

## Summary
`hybrid_dispatch.cuh` calls `gin.red_add_rel` to update the `scaleout_channel_signaled_tail` slot. On the local-rank bypass path this dispatches to `ptx::red_add_rel_sys` in `handle.cuh`, emitting `red.release.sys.global.add.u64`.
After tracing through `get_sym_ptr`, we are wondering whether `.gpu` might also be sufficient here, or whether there is a specific reason `.sys` is required that we are missing?

## Question about the existing comment
the comment`NOTES: the "release" scope will be `sys` for the local rank (we may involve NVLink so not gpu)` mentions NVLink, but the only NVLink-bearing path goes through ncclTeamTagLsa/ncclTeamTagWorld, not ncclTeamTagRail. Is there a code path we missed where Rail can land on an NVLink peer?

## Proposed change
Add a red_add_rel_gpu overload (or a scope tag) to NCCLGin so the call site can opt into .gpu release:
```
// handle.cuh — new overload alongside existing red_add_rel
template
__device__ __forceinline__
void red_add_rel_gpu(dtype_t* sym_ptr, const dtype_t& value,
const int& dst_rank_idx,
const int& extra_options = 0) const {
const auto dst_ptr = get_sym_ptr(sym_ptr, dst_rank_idx);
if (dst_ptr != nullptr) {
ptx::red_add_rel_gpu(dst_ptr, value); // <-- .gpu instead of
.sys
} else {
// RDMA path unchanged: atomic provides system-scope ordering
gin.signal(...);
}
}
```
The same change applies to hybrid_combine.cuh by the same reasoning.
## Questions
1. Is there a non-local consumer of scaleout_channel_signaled_tail_ptr we might have missed ?
2. Is the NVLink mention in the comment a leftover from an earlier design, or does Rail ever land on an NVLink peer in some configuration?

### update
1. ~30% improvement when ep16 with 12 SMs
2. 5% improvement when ep16 with 16 SMs