AnsweredAssumed Answered

Irreducible ControlFlow Detected

Question asked by petr.machacek on Jun 11, 2013
Latest reply on Jun 12, 2013 by himanshu.gautam

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[S];

  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[n])

       {

            lab[n] = label;    

            *m = true;

       }

       else

       {

            label = lab[n];

       }

  }

}

Outcomes