Deterministic mode silently corrupts scratch-dependent counter reservations

[eric-heiden]

Bug Description

Deterministic mode can silently produce incorrect results when a consumed-return counter atomic is reached through control flow that depends on a global scratch write earlier in the same kernel.

Warp's deterministic counting pass suppresses the scratch write. The counting and execution passes can therefore encounter different counter-reservation sequences. Warp currently accepts this divergence and replays missing/incorrect reservation records instead of raising an error.

Even if this pattern remains unsupported, it should fail clearly rather than return plausible but incorrect output.

Minimal Reproducer

import warp as wp


wp.config.deterministic = wp.DeterministicMode.RUN_TO_RUN
wp.config.deterministic_max_records = 1


@wp.kernel
def reserve_after_scratch_write(
    values: wp.array(dtype=wp.int32),
    scratch: wp.array(dtype=wp.int32),
    count: wp.array(dtype=wp.int32),
    slots: wp.array(dtype=wp.int32),
):
    tid = wp.tid()
    scratch[tid] = values[tid]

    if scratch[tid] != 0:
        slot = wp.atomic_add(count, 0, 1)
        slots[slot] = tid


thread_count = 4
values = wp.ones(thread_count, dtype=wp.int32, device="cuda:0")
scratch = wp.zeros(thread_count, dtype=wp.int32, device="cuda:0")
count = wp.zeros(1, dtype=wp.int32, device="cuda:0")
slots = wp.full(thread_count, -1, dtype=wp.int32, device="cuda:0")

wp.launch(
    reserve_after_scratch_write,
    dim=thread_count,
    inputs=[values, scratch, count, slots],
    device="cuda:0",
)

print("scratch:", scratch.numpy())
print("count:", count.numpy())
print("slots:", slots.numpy())

With NOT_GUARANTEED:

scratch: [1 1 1 1]
count: [4]
slots: [0 1 2 3]

With either RUN_TO_RUN or GPU_TO_GPU on the tested system:

scratch: [1 1 1 1]
count: [0]
slots: [ 0 -1 -1 -1]

No warning or exception is raised. Increasing deterministic_max_records does not fix the semantic mismatch.

Real-world Trigger

This was found in MuJoCo-Warp 3.11.0 convex collision handling. Its CCD kernel reserves an EPA scratch slot with one consumed-return atomic_add, writes and reads global scratch data keyed by that slot, and later reserves contact output with another consumed-return atomic_add.

The suppressed scratch writes make the later reservation/control-flow sequence differ between passes. Convex contacts can disappear or change, causing resting bodies to fall through or tip off surfaces. Primitive collision pairs in the same scene are unaffected, and CPU MuJoCo agrees with the non-deterministic Warp result.

Expected Behavior

Warp should detect that the deterministic replay does not match the counting pass and raise an actionable error, for example:

Deterministic counter replay diverged because slot allocation depends on side
effects suppressed during the counting pass. Split computation and allocation
into separate kernels.

One possible validation is to check, per thread, that phase 1 consumes exactly the records produced by phase 0 and that each replayed record matches the counter target/index and reservation value. This needs both an immediate check for extra or mismatched reservations and an end-of-pass check for missing reservations.

System Information

• Warp: 1.16.0
• MuJoCo-Warp: 3.11.0
• CUDA Toolkit: 12.9
• Driver API: 13.0
• GPU: NVIDIA GeForce RTX 4090 Laptop GPU (sm_89)
• OS: Windows