[BUG]: Atomic_ref.fetch_add "segfaults" on floating point local shared memory value in cluster launch

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

libcu++

### Describe the bug

The following kernel fails when launched as a cluster:
```
template
__global__ void test() {
__shared__ T x;

cuda::atomic_ref atomic_data(x);
atomic_data.fetch_add(T(1), cuda::memory_order_relaxed);
}
```
When it is not launched as a cluster, it works fine.

### How to Reproduce

Run `bash reproducer.cu`

reproducer.cu:

```
#if 0
set -ex
nvcc -run -arch sm_90 $0
exit 0
#endif

#include
#include
#include

// Macro for checking cuda errors following a cuda launch or api call
#define CUDA_CHECK(ans) \
{ \
gpuAssert((ans), __FILE__, __LINE__); \
}

inline void gpuAssert(cudaError_t code, const char *file, int line,
bool abort = true)
{
if (code != cudaSuccess)
{
fprintf(stderr, "CUDA error: %s at %s:%d\n", cudaGetErrorString(code), file,
line);
if (abort)
exit(code);
}
}

template
inline __device__ T* launder_pointer(T* ptr) {
T * unknown_ptr;
asm("mov.b64 %0, %1;": "=l"(unknown_ptr): "l"(ptr));
return unknown_ptr;
}

template
__global__ void test() {
__shared__ T x;

cuda::atomic_ref atomic_data(x);
atomic_data.fetch_add(T(1), cuda::memory_order_relaxed);
}

struct cluster_launch{};
struct normal_launch{};

template
void launch(launch_type lt) {
fprintf(stderr, "%-100s ", __PRETTY_FUNCTION__);

if (std::is_same_v) {
test<<>>();
} else {
cudaLaunchConfig_t config = {0};
// The grid dimension is not affected by cluster launch, and is still enumerated
// using number of blocks.
// The grid dimension should be a multiple of cluster size.
config.gridDim = cluster_dims;
config.blockDim = 1;

cudaLaunchAttribute attribute[1];
attribute[0].id = cudaLaunchAttributeClusterDimension;
attribute[0].val.clusterDim.x = cluster_dims; // Cluster size in X-dimension
attribute[0].val.clusterDim.y = 1;
attribute[0].val.clusterDim.z = 1;
config.attrs = attribute;
config.numAttrs = 1;
CUDA_CHECK(cudaLaunchKernelEx(&config, test));
}
CUDA_CHECK(cudaDeviceSynchronize());
fprintf(stderr, "SUCCESS\n");
}

int main() {
CUDA_CHECK(cudaDeviceSynchronize());
launch(normal_launch{});
launch(normal_launch{});
launch(normal_launch{});
launch(normal_launch{});

launch(cluster_launch{});
launch(cluster_launch{});
launch(cluster_launch{});
launch(cluster_launch{});

CUDA_CHECK(cudaDeviceSynchronize());
}

```

Output:
```
bash bug_atomic_ref.cu
+ nvcc -run -arch sm_90 bug_atomic_ref.cu
void launch(launch_type) [with T = int; int cluster_dims = 1; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 2; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 1; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 2; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 1; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 2; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 1; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 2; launch_type = cluster_launch] CUDA error: an illegal instruction was encountered at bug_atomic_ref.cu:74
```
Expected output:
```
bash bug_atomic_ref.cu
+ nvcc -run -arch sm_90 bug_atomic_ref.cu
void launch(launch_type) [with T = int; int cluster_dims = 1; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 2; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 1; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 2; launch_type = normal_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 1; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = int; int cluster_dims = 2; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 1; launch_type = cluster_launch] SUCCESS
void launch(launch_type) [with T = float; int cluster_dims = 2; launch_type = cluster_launch] SUCCESS
```

### Expected behavior

The kernel should work in normal launch and cluster launch.

### Reproduction link

_No response_

### Operating System

NA

### nvidia-smi output

NA

### NVCC version

all.