[BUG]: `cooperative_groups::thread_block::sync()` produces incorrect results on SM_120 with CUDA 13.2

### Is this a duplicate?

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

### Type of Bug

Runtime Error

### Component

Not sure

### Describe the bug

When testing the cudf CUDA 13.2 build on `sm_120`, we observed groupby failures that do not appear on other GPUs or with other CUDA Toolkit versions i.e. https://github.com/rapidsai/cudf/issues/22159

The fix is simple but odd: simply replacing this line https://github.com/rapidsai/cudf/blob/601bbcf9a7fd8d055f8b7fb34cac9c79c7fd4173/cpp/src/groupby/hash/compute_shared_memory_aggs.cu#L233 with `__syncthreads`. (see https://github.com/rapidsai/cudf/pull/22152)

It seems there is an issue with `cooperative_groups::thread_block::sync()`, which is expected to be semantically equivalent to `__syncthreads()`, but does not behave identically in practice. By inspecting the generated SASS, we observed that both `thread_block::sync()` and `__syncthreads()` use `WARPSYNC.COLLECTIVE.ALL`. However, `__syncthreads()` additionally emits `WARPSYNC.ALL`, which is absent in the cooperative groups version.

Based on offline discussions with @pciolkosz, it's suspicious that `WARPSYNC` is executed before `ENDCOLLECTIVE` in the block sync SASS, since collective blocks are supposed to keep participating threads converged by construction and explicitly disallow `WARPSYNC` inside them.

For reference, I have attached the SASS outputs from both cases: [sass_comp.zip](https://github.com/user-attachments/files/26905369/sass_comp.zip)

### How to Reproduce

Build cudf https://github.com/rapidsai/cudf/tree/2ea649d59a44937b04bb3ae3f21d440fcf58ddd6 with the 13.2 devcontainer and run GROUPBY unit tests on `sm_120`.

### Expected behavior

Using `block.sync()` should pass all GROUPBY unit tests.