Skip to content

Chained and in kernels does not short-circuit — nullable array access evaluated unconditionally #1329

Closed

Description

@eric-heiden

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.

Metadata

Metadata

Assignees

Labels

bugSomething isn't working

Type

Projects

No projects

    Milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions