Hello,
I'm trying to compile my OpenCL kernel but it allways finishes with this error:
Error:E010:Irreducible ControlFlow Detected
Can someone explain to me, where is the error in my code?
Maybe it is a bug in Stream compiler (http://devgurus.amd.com/message/1127626#1127626) but may be it is my fault...
Here is the code of the kernel:
#define CONNECTED(a,b) (((a>0) && (b>0))||((a==0) && (b==0)))
__kernel void myKernel(
int I,
__global uchar *input,
int stride,
__global uint *lab,
__global bool *m,
int pixCount,
int width
)
{
int id = (int)get_global_id(0);
int H = pixCount / width;
int S=0, E=0, step=0;
switch(I)
{
case 0:
if(id>=width)
{
return;
}
S = id;
E = width * (H - 1) + id;
step = width;
break;
case 1:
if(id >= H)
{
return;
}
S = id * width;
E = S + width - 1;
step = 1;
break;
case 2:
if(id >= width)
{
return;
}
S = width * (H - 1) + id;
E = id;
step = -width;
break;
case 3:
if(id >= H)
{
return;
}
S = (id + 1) * width - 1;
E = id * width;
step = -1;
break;
}
uint label = lab;
for(int n = S + step; n != E + step; n += step)
{
uint r1 = n / width;
uint c1 = n % width;
uint r2 = (n-step) / width;
uint c2 = (n-step) % width;
if( CONNECTED(input[r1*stride + c1], input[r2*stride + c2]) && label < lab
{
lab
*m = true;
}
else
{
label = lab
}
}
}
I have just removed the compilation error by editing my code (red highlighted). It seems to be an compiler error, so I'll send it to streamdeveloper@amd.com.
#define CONNECTED(a,b) (((a>0) && (b>0))||((a==0) && (b==0)))
void avoidErrorFunc(uchar i1, uchar i2, uint *label, __global uint *labn, __global bool *m)
{
if( CONNECTED(i1, i2) && *label < *labn)
{
*labn = *label;
*m = true;
}
else
{
*label = *labn;
}
}
__kernel void myKernel(
int I,
__global uchar *input,
int stride,
__global uint *lab,
__global bool *m,
int pixCount,
int width
)
{
int id = (int)get_global_id(0);
int H = pixCount / width;
int S=0, E=0, step=0;
switch(I)
{
case 0:
if(id>=width)
{
return;
}
S = id;
E = width * (H - 1) + id;
step = width;
break;
case 1:
if(id >= H)
{
return;
}
S = id * width;
E = S + width - 1;
step = 1;
break;
case 2:
if(id >= width)
{
return;
}
S = width * (H - 1) + id;
E = id;
step = -width;
break;
case 3:
if(id >= H)
{
return;
}
S = (id + 1) * width - 1;
E = id * width;
step = -1;
break;
}
uint label = lab;
for(int n = S + step; n != E + step; n += step)
{
uint r1 = n / width;
uint c1 = n % width;
uint r2 = (n-step) / width;
uint c2 = (n-step) % width;
avoidErrorFunc(input[r1*stride + c1], input[r2*stride + c2], &label, &lab
}
}
What driver, GPU and OS? I do not see compile error in kernel analyzer , with catalyst 13.6 linux when SI & NI cards are selected.
Here are the informations:
OS: Windows 7 Home Premium (64bit)
GPUs: HD5670 and HD5750
APP SDK: 12.8 Windows-64
Driver: OpenCL 1.2 AMD-APP (1124.2)
Catalyst: 13.4
My KernelAnalyzer (v 1.12.1288) displays the error as well.
Hi,
I also see the same with 13.6 driver, HD 7870, Win7-64. I am forwarding this to relevant people.