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

[minansys]

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:

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:

./enzyme_cuda_cache_repro

Observed output:

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:

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:

%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.

[minansys]

@wsmoses Here is the issue I mentioned yersterday