[BUG] Kernel if short circuiting

[ped02]

Bug Description

Prior to warp v1.5 (tested with v1.4.0), the condition if reset and reset[tid] is able to run when reset is None. After upgrading to v1.5, the kernel cause Warp CUDA Error 700: an illegal memory access was encountered. Not sure if upgrading was supposed to report illegal memory access or causing illegal memory access.

import warp as wp

@wp.kernel
def test_kernel(
    reset: wp.array(dtype=bool),
    # outputs
    buffer: wp.array(dtype=int)
):

    tid = wp.tid()

    ## Works only before v1.5
    if reset and reset[tid]:
        buffer[tid] = tid

    ## Solution for > v1.5
    # if reset:
    #     if reset[tid]:
    #         buffer[tid] = tid


def main():

    reset = wp.array([True, True, False, True, False, False])

    buffer = wp.empty(reset.shape, dtype=int)

    wp.launch(
        test_kernel,
        dim=buffer.shape,
        inputs = [
            # reset
            None
        ],
        outputs = [
            buffer
        ]
    )

    print(f'{buffer.numpy()=}')

if __name__ == '__main__':
    main()

System Information
CUDA Version: 12.5
OS Ubuntu 22.04
Python Version: 3.10.12

System Information

No response

[Hantao-Ye]

As they mentioned in the documentation, short-curcuit evaluation is not supported in kernels and user functions.

[shi-eric]

A git bisect indicates that the behavior change started in 3f31e2c#diff-a7be079c872d39bb82545a63383217ea48bc14d5d708f0a99b025d55f039f1a4 but it's unclear to us why the CUDA Error 700 starts getting reported.

@Hantao-Ye is correct that we don't support short-circuit evaluation right now, so in v1.4.0 and in v1.5.0 we're loading from a malformed address:

        var_0 = builtin_tid1d();
        // if reset and reset[tid]:                                                               <L 13>
        var_1 = wp::address(var_reset, var_0);
        var_2 = wp::load(var_1);
        var_3 = var_reset && var_2;
        if (var_3) {
            // buffer[tid] = tid                                                                  <L 14>
            wp::array_store(var_buffer, var_0, var_0);
        }

It's best to use the solution you identified.