5 Replies Latest reply on Jun 12, 2013 12:10 AM by himanshu.gautam

    Irreducible ControlFlow Detected

    petr.machacek

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

             }

        }

      }

        • Re: Irreducible ControlFlow Detected
          petr.machacek

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

           

                avoidErrorFunc(input[r1*stride + c1], input[r2*stride + c2], &label, &lab[n], m);

            }

          }