[FEA]: Use mbarriers to improve barrier/communication latency in cluster segmented TopK

### Is this a duplicate?

- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

CUB

### Is your feature request related to a problem? Please describe.

Cluster barriers are quite expensive because they need to do an L2 round-trip (CTAs could be communicating through global memory). Furthermore they need all CTA's in a cluster to stay around until the end. Finally communication through DSMEM followed by a barrier is generally less efficient than combining the data and the information about how much data is sent into one instruction.

### Describe the solution you'd like

One can relax the barrier arrival but then needs additional fences to "publish" e.g. smem state to the other CTAs in the same cluster. There is one special fence that allows publishing mbarrier-initialization without the L2 roundtrip. So the natural solution is to have a single relaxed cluster barrier at the start of the kernel to publish mbarriers and then only use those in combination with async atomics/stores that can add to the transaction (tx) count on the remote mbarrier.

This issue is done when the described approach is implemented, tested and benchmarked to be faster than the current approach.

### Describe alternatives you've considered

_No response_

### Additional context

_No response_