[BUG]: Invalid codegen when combining 256B LD/ST with L2 access policies

### 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

Compile-time Error

### Component

libcu++

### Describe the bug

With the Blackwell architecture we introduced 256 byte vectorized load/store instructions: `ld/st.global.v8` aka `LDG/STG.256`. Combining these instructions with L2 access policies (`cuda::annotated_ptr`or `cuda::associate_access_policy`) leads to a compiler error "illegal instruction".

### How to Reproduce

https://godbolt.org/z/5P8j43PrG -> Illegal instruction (core dumped)

If we switch to 128B LD/ST instead, the cache hints get applied in PTX (`ld/st.global.L2::cache_hint.v4`) but are somehow dropped when lowered to SASS.

### Expected behavior

https://godbolt.org/z/7v3Pdjshj

Lowers to `LDG/STG.E.ELL2.256` where `ELL2` means "evict last L2" - as expected.

### Reproduction link

https://godbolt.org/z/5P8j43PrG

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

v13.0.1 on Godbolt