Repository navigation
Chained and in kernels does not short-circuit â nullable array access evaluated unconditionally #1329
Description
Summary
Warp's codegen for chained and expressions does not short-circuit. When a kernel guards an array element access behind a nullable-array truthiness check using a single if with and, the array access is evaluated unconditionally â even when the array is None/null.
This affects both CPU and CUDA backends and appears to be a codegen issue rather than a runtime issue.
Expected behavior
if arr and idx >= 0 and arr[idx] == 0:should stop evaluating after arr is falsy (null), consistent with Python semantics. The element access arr[idx] should never execute when arr is None.
Actual behavior
- CUDA:
arr[idx]is evaluated with a null pointer â CUDA error 700 (illegal memory access), poisons the entire CUDA context. - CPU:
arr[idx]is evaluated with a null pointer â segfault or garbage read (undefined behavior).
Workaround
Nested if statements work correctly:
if arr:
if idx >= 0 and arr[idx] == 0:
...Minimal reproducer
"""
Minimal reproducer: Warp codegen does not short-circuit chained `and`
when the first operand is a nullable array.
The array element access in the third operand is evaluated even when
the first operand (the array itself) is None, causing:
- CUDA: error 700 (illegal memory access)
- CPU: segfault / undefined behavior
Expected: `if arr and idx >= 0 and arr[idx] == 0:` should skip
arr[idx] when arr is None.
Actual: arr[idx] is evaluated unconditionally.
Workaround: use nested `if` statements.
"""
import warp as wp
wp.init()
@wp.kernel
def chained_and_kernel(
arr: wp.array(dtype=int),
result: wp.array(dtype=int),
):
tid = wp.tid()
# BUG: arr[tid] is evaluated even when arr is None
if arr and tid >= 0 and arr[tid] == 0:
result[tid] = -1
return
result[tid] = 1
@wp.kernel
def nested_if_kernel(
arr: wp.array(dtype=int),
result: wp.array(dtype=int),
):
tid = wp.tid()
# WORKAROUND: nested if correctly guards the access
if arr:
if tid >= 0 and arr[tid] == 0:
result[tid] = -1
return
result[tid] = 1
for device in ("cpu", "cuda:0"):
if device.startswith("cuda") and not wp.is_cuda_available():
print(f"Skipping {device} (no CUDA)")
continue
print(f"\n=== Device: {device} ===")
result = wp.zeros(3, dtype=int, device=device)
# Nested if with None array â should work fine
print("Nested if (None array)... ", end="", flush=True)
try:
wp.launch(nested_if_kernel, dim=3, inputs=[None, result], device=device)
wp.synchronize_device(device)
print(f"OK result={result.numpy().tolist()}")
except Exception as e:
print(f"FAIL {e}")
# Chained and with None array â triggers the bug
print("Chained and (None array)... ", end="", flush=True)
try:
wp.launch(chained_and_kernel, dim=3, inputs=[None, result], device=device)
wp.synchronize_device(device)
print(f"OK result={result.numpy().tolist()}")
except Exception as e:
print(f"FAIL {e}")Environment
- Warp: 1.7.2 (also reproduced on 1.6.x)
- Python: 3.12
- OS: Ubuntu 22.04
- GPU: NVIDIA RTX (CUDA 12.x) â but bug also reproduces on CPU backend
Impact
Any kernel that uses a nullable array parameter and guards element access with a chained and expression will hit undefined behavior. The nested if workaround is straightforward but the behavior is surprising since it silently diverges from Python's short-circuit semantics.
Activity
c0d1f1ed commentedon Mar 30, 2026
Implemented in 5d81609.
Fix short-circuit evaluation of chained `and`/`or` in kernels (NVIDIAâĶ
Update docs to reflect short-circuit evaluation support (NVIDIAGH-1329)
Add changelog entry for short-circuit evaluation fix (NVIDIAGH-1329)
Merge branch 'ncapens/fix-short-circuiting' into 'main'