cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

Irreducible ControlFlow Detected

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 = label;    

            *m = true;

       }

       else

       {

            label = lab;

       }

  }

}

0 Likes
5 Replies

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, m);

  }

}

0 Likes

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.

0 Likes

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

0 Likes

My KernelAnalyzer (v 1.12.1288) displays the error as well.

0 Likes

Hi,

I also see the same with 13.6 driver, HD 7870, Win7-64. I am forwarding this to relevant people.

0 Likes