[BUG]: cuda::std::memcpy asserts "source range is invalid" for shared pointer in debug build

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

`cuda::std::memcpy` / `cuda::std::memset` fires a spurious "source range is invalid" assertion in debug builds (CCCL_ENABLE_ASSERTIONS, no -DNDEBUG) on NVCC 13+ when the source / destination pointer is in shared memory but the compiler cannot prove this statically.

CCCL's `cuda::std::memcpy` debug path calls `__is_valid_address_range(src, n)` which calls `isspacep.shared(src) → TRUE`, then calls `__is_smem_valid_address_range(src, n)`
Inside that function, after `_CCCL_ASSUME(__p)` is applied following the first `isspacep.shared` result, NVCC 13+ generates INVERTED PTX logic for the range validation: the `assertfail` branch fires when the smem range IS valid, rather than when it is NOT. This causes a spurious assertion for any shared-memory pointer when the compiler cannot prove address-space membership statically (i.e. it emits `isspacep.shared` rather than `cvta.shared.u64` to identify the address space at runtime).

CCCL files involved:
- `libcudacxx/include/cuda/std/__cstring/memcpy.h` (assertion at line 49)
- `libcudacxx/include/cuda/std/__cstring/memset.h` (assertion at line 38)
- `libcudacxx/include/cuda/__memory/check_address.h (__is_smem_valid_address_range)`
- `libcudacxx/include/cuda/__memory/address_space.h (_CCCL_ASSUME misuse)`

### How to Reproduce

I was not able to reproduce in a goldbolt session but the following reproducer should show the issue

```cpp
// Build (debug, assertions enabled, no -DNDEBUG):
// nvcc -g -O0 -std=c++20 -DCCCL_ENABLE_ASSERTIONS \
// -I/libcudacxx/include \
// -o repro1 repro1.cu && ./repro1
//

#include
#include
#include
#include

static constexpr int BLOCK = 128;
static constexpr int BUF_SIZE = 1024;

// ------------------------------------------------------------------
// Kernel A (reproducer): smem pointer made opaque to the compiler
// via an empty inline-asm constraint so the compiler cannot trace
// it back to the __shared__ declaration at compile time. This forces
// NVCC to emit isspacep.shared in the CCCL assertion check, where the
// NVCC 13.1 code-generation bug triggers the spurious assertfail.
//
// PTX emitted: isspacep.shared %p5, %rd7 (confirmed via -ptx)
// Expected: FAILS with "memcpy: source range is invalid"
// ------------------------------------------------------------------
__global__ void kernel_A_opaque_smem(uint8_t const* d_src)
{
__shared__ __align__(16) uint8_t smem_buf[BUF_SIZE];
for (int i = threadIdx.x; i < BUF_SIZE; i += blockDim.x) smem_buf[i] = d_src[i];
__syncthreads();
if (threadIdx.x == 0) {
// Inline asm makes the compiler treat ptr as an unknown generic pointer —
// the same situation that arises in libcudf when the __shared__ pointer passes
// through function parameters, cuda::std::assume_aligned, and
// struct-member store/loads (e.g. rle_stream.cuh s_start).
uint8_t* ptr = smem_buf;
asm volatile("" : "+l"(ptr) : :); // ptr origin now opaque to compiler
int32_t val;
cuda::std::memcpy(&val, ptr + 4, sizeof(val));
printf("A val=0x%08x\n", val);
}
}

int main(int argc, char* argv[])
{
uint8_t* d_src;
cudaMalloc(&d_src, BUF_SIZE);
cudaMemset(d_src, 0x42, BUF_SIZE);
kernel_A_opaque_smem<<<1, BLOCK>>>(d_src);
cudaError_t e = cudaDeviceSynchronize();
cudaFree(d_src);
if (e != cudaSuccess) {
fprintf(stderr, "FAILED: %s\n", cudaGetErrorString(e));
return 1;
}
printf("PASSED\n");
return 0;
}

```
Compile command:
```
nvcc -g -O0 -std=c++20 -DCCCL_ENABLE_ASSERTIONS -o repro1 repro1.cu
```

### Expected behavior

No assert. Release build (without asserts) shows no ill effects.

### Reproduction link

_No response_

### Operating System

Ubuntu Linux 22.04

### nvidia-smi output

```
Wed Aug 19 15:29:30 2026
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 610.57.04 KMD Version: 610.57.04 CUDA UMD Version: 13.3 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA T400 4GB On | 00000000:01:00.0 On | N/A |
| 38% 39C P8 N/A / 31W | 1227MiB / 4096MiB | 5% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
| 1 NVIDIA RTX A6000 On | 00000000:C1:00.0 Off | Off |
| 30% 37C P8 21W / 300W | 15MiB / 49140MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| 0 N/A N/A 3803 G /usr/lib/xorg/Xorg 416MiB |
| 0 N/A N/A 3992 G /usr/bin/gnome-shell 169MiB |
| 0 N/A N/A 18391 G ...rack-uuid=3190708988185955192 354MiB |
| 0 N/A N/A 96181 G ...rack-uuid=3190708988185955192 205MiB |
| 1 N/A N/A 3803 G /usr/lib/xorg/Xorg 4MiB |
+-----------------------------------------------------------------------------------------+
```

### NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2026 NVIDIA Corporation
Built on Tue_Jun_09_02:43:40_PM_PDT_2026
Cuda compilation tools, release 13.3, V13.3.73
Build cuda_13.3.r13.3/compiler.38244171_0

Start with libcudacxx/include/cuda/__memory/check_address.h and __is_smem_valid_address_range, then inspect the related _CCCL_ASSUME use in address_space.h. Review the assertion call sites in __cstring/memcpy.h and __cstring/memset.h, and build the supplied CUDA reproducer with CCCL_ENABLE_ASSERTIONS. Done means the opaque shared-memory pointer no longer triggers a spurious assertion in the debug build.

Written by the indexing model from the issue text.