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