cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

davibu
Journeyman III

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

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.

 

0 Likes
3 Replies
himanshu_gautam
Grandmaster

hi davibu,

i tried the problem you specified on vista X64.It indeed show only a warning in the first case and do compile.But i do not think declaring a label inside a case

causes any accuracy issue.May be AMD doesn't consider the problem critical enough to be reported as an error.A warning is definitely noticeable.

0 Likes
genaganna
Journeyman III

Originally posted 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.

 

 



Thanks for reporting first issue. Upcoming releases will give appropriate error for this.

Could you please tell me where buggy code gererated(You second case)? Please give kernel code which shows this issue.

0 Likes

Originally posted by: genaganna

 

Thanks for reporting first issue. Upcoming releases will give appropriate error for this.

 

Could you please tell me where buggy code gererated(You second case)? Please give kernel code which shows this issue.

 

 

I will try to pick the exact version from the Mercurial repository: http://src.luxrender.net/luxrays

I was experiencing very long compilation time (i.e. up to 300 seconds, several GBs of memory used, etc.) for a kernel too. I will look if I can report both problems in a single shot.

 

0 Likes