NVIDIA / warp

A Python framework for high performance GPU simulation and graphics
https://nvidia.github.io/warp/
Other
4.23k stars 242 forks source link

[BUG] Incorrect transformations of impossible code #299

Closed steinraf closed 1 month ago

steinraf commented 2 months ago

Bug Description

In some circumstances its possible to create code which should raise an error when run in python, but gets transformed in a way that does not produce an error in the autogenerated cuda kernel code.

If a variable is initialized in a branch that is not taken, no error is raised and the value gets initialized anyways

import warp as wp

@wp.kernel
def test(x: wp.array(dtype=float)):
    if False:
        a = 1.0

    x[0] = a

x = wp.zeros(1)
wp.launch(test, dim=1, inputs=[x])

wp.synchronize()

print(x)
# > [1.]

The kernel is translated to the following:

const bool var_0 = false;
const wp::float32 var_1 = 1.0;
const wp::int32 var_2 = 0;
//---------
// forward
// def test(x: wp.array(dtype=float)):                                                    <L 5>
// if False:                                                                              <L 6>
if (var_0) {
    // a = 1.0                                                                            <L 7>
}
// x[0] = a                                                                               <L 11>
wp::array_store(var_x, var_2, var_1);

Expected behaviour: UnboundLocalError: local variable 'a' referenced before assignment or error: identifier "a" is undefined

Another, arguably less relevant but similar, example is when computing divisions by zero with known constants. In python this usually causes a runtime error and in C/C++ it is undefined behaviour and creates a warning if the constant is known to be 0 at compile time.

@wp.kernel
def test(x: wp.array(dtype=float)):
    d = 0.0
    x[0] = x[0] / d

x = wp.ones(1)
wp.launch(test, dim=1, inputs=[x])

wp.synchronize()

print(x)

# > [inf]

This produces the following kernel code

const wp::float32 var_0 = 0.0;
const wp::int32 var_1 = 0;
wp::float32* var_2;
wp::float32 var_3;
wp::float32 var_4;
//---------
// forward
// def test(x: wp.array(dtype=float)):                                                    <L 47>
// d = 0.0                                                                                <L 48>
// x[0] = x[0] / d                                                                        <L 49>
var_2 = wp::address(var_x, var_1);
var_3 = wp::load(var_2);
var_4 = wp::div(var_3, var_0);
wp::array_store(var_x, var_1, var_4);

Expected behaviour: ZeroDivisionError: float division by zero or warning #39-D: division by zero

System Information

Warp 1.3.0 initialized: CUDA Toolkit 12.5, Driver 12.2

christophercrouzet commented 2 months ago

Hi @steinraf, thank you for reporting this issue!

Regarding the division by 0, you could try setting wp.config.verify_fp = True, which should catch such issues and log them.

As for accessing that scoped variable, when the if condition can be evaluated to False at compile time, as it is the case in your example, we should be able to discard the dead code altogether, which is something that @eric-heiden's current work around static expressions might help addressing.

However, if the if condition requires to be evaluated at runtime, then it's tricky. The issue is that Python has different scoping rules than the C++/CUDA languages that Warp is transpiling that Python code to, so it's not really possible to nicely map between both in this case.

steinraf commented 1 month ago

Yes, the floating point verification catches that error during runtime, I was wondering more if there was a way to catch these errors (or other ones) by getting a warning/error during compilation. Is there any way to access the compilation logs in general?

The work on static expressions sounds exciting, looking forward to that!

I agree, the both languages don't interact nicely regarding that. At this point it would be most helpful for me if I looked at the code generation to see how these cases are handled. Until then its probably best just not to initialize variables in branches.

christophercrouzet commented 1 month ago

The most logs that Warp can currently output can be enabled with the following configuration values:

wp.config.mode = "debug"
wp.config.verbose = True
wp.config.verify_fp = True
wp.config.verify_cuda = True

That includes compilation warnings/errors, if any.