davibu

[BUG] Compiler problem with switch() statement and variable declaration

Discussion created by davibu on Sep 1, 2010
Latest reply on Sep 3, 2010 by davibu

I have experience this problem with a large complex kernel and I try to exemplify the problem here with an example.

The AMD OpenCL C compiler compiles the following code:

__kernel void test(
        __global int *v,
        ) {
    int gid = get_global_id(0);

    switch(gid) {
        case 0:
            int a = 0; // <=== Problem here, missing { ... }
            gid = a + 1;
            break;
    }

    v[gid] = gid;

}

by firing only a warning (i.e. "warning: a declaration cannot have a label"). First of all, the compiler should refuse to compile the above code because it has an error. The right code is:

__kernel void test(
        __global int *v,
        ) {
    int gid = get_global_id(0);

    switch(gid) {
        case 0:

{ // <=== Addition
            int a = 0;
            gid = a + 1;

            break;

} // <=== Addition
    }

    v[gid] = gid;

}

Second, the compiler seems to generate some buggy code like if multiple variables were sharing the same register at the same time. This lead indeed to some very wired behaviour of the code.

The problem was found on Linux Ubuntu 10.04 with SDK 2.2 and Catalyst 10.8.

Both NVIDIA OpenCL and Apple OpenCL show the right behaviour: they refuse to compile the code with an error.

 

Outcomes