CUDA reverse cache malloc is marked nonnull and can cause illegal memory access

# CUDA reverse cache malloc is marked nonnull and can cause illegal memory access

## Summary

Enzyme-generated CUDA reverse code can emit device-side cache allocations as `nonnull` and immediately use the returned pointer without checking for allocation failure. With CUDA's default device malloc heap (`8388608` bytes on my run), this reproducer reports `cudaErrorIllegalAddress` at `cudaDeviceSynchronize()`.

The reproducer is generic: it differentiates a CUDA device function with a 100-iteration fixed-point loop and loop-carried values. It does not depend on any application code.

## Environment

- Enzyme package: `0.0.261`
- Clang: `19.1.7`
- CUDA: `12.8`
- GPU arch used: `sm_80`
- Relevant Enzyme flags:
- `-mllvm -enzyme-loose-types=1`
- `-mllvm -enzyme-detect-readthrow=0`
- `-mllvm -enzyme-coalese=1`
- `-mllvm -enzyme-noalias=1`
- `-mllvm -enzyme-phi-restructure=1`
- also reproduced with `-mllvm -enzyme-max-cache=1`

## Reproducer

Save as `enzyme_cuda_cache_repro.cu`:

https://tinyurl.com/274884zq

## Build and Run

Build without `-enzyme-max-cache=1`:

```bash
clang++ -O3 -x cuda --cuda-gpu-arch=sm_80 --cuda-path=/path/to/cuda \
-fplugin=/path/to/libClangEnzyme.so \
-mllvm -enzyme-loose-types=1 -mllvm -enzyme-detect-readthrow=0 \
-mllvm -enzyme-coalese=1 -mllvm -enzyme-noalias=1 \
-mllvm -enzyme-phi-restructure=1 \
enzyme_cuda_cache_repro.cu \
-L/path/to/cuda/lib64 -Wl,-rpath,/path/to/cuda/lib64 -lcudart -ldl -lrt -pthread \
-o enzyme_cuda_cache_repro
```

Run:

```bash
./enzyme_cuda_cache_repro
```

Observed output:

```text
n=32768 heap_bytes=8388608 heap_mode=default
launch: no error
sync: an illegal memory access was encountered
reproduced Enzyme-generated device cache allocation failure
```

I also reproduced the same default-heap output when adding `-mllvm -enzyme-max-cache=1`.

## Generated IR Evidence

Generating device LLVM IR with:

```bash
clang++ --cuda-device-only -S -emit-llvm -O3 -x cuda \
--cuda-gpu-arch=sm_80 --cuda-path=/path/to/cuda \
-fplugin=/path/to/libClangEnzyme.so \
-mllvm -enzyme-loose-types=1 -mllvm -enzyme-detect-readthrow=0 \
-mllvm -enzyme-coalese=1 -mllvm -enzyme-noalias=1 \
-mllvm -enzyme-phi-restructure=1 \
enzyme_cuda_cache_repro.cu -o repro.ll
```

shows repeated cache allocations like:

```llvm
%88 = tail call noalias nonnull ptr @malloc(i64 %84)
%90 = getelementptr inbounds i8, ptr %88, i64 %87
tail call void @llvm.memset.p0.i64(ptr nonnull align 1 %90, i8 0, i64 %89, i1 false)
tail call void @llvm.memcpy.p0.p0.i64(ptr nonnull align 1 %88, ptr align 1 %69, i64 %87, i1 false)
```

There is no check for `malloc == null`. The pointer is marked `nonnull`, and subsequent `memset`/`memcpy`/stores also use `nonnull`. On CUDA, device `malloc` can fail when the device heap is exhausted or fragmented; that becomes an illegal memory access in the generated reverse kernel.

## Expected Behavior

For CUDA/device-side generated cache allocation, Enzyme should not assume `malloc` is non-null unless that is guaranteed. Possible expected behavior:

- Avoid emitting `nonnull` on device malloc cache allocations.
- Guard generated uses against null allocation failure.
- Provide an option/runtime path that reports device cache allocation failure instead of creating an illegal memory access.