11 Replies Latest reply on Aug 20, 2010 6:27 PM by ryta1203

    Irreducible ControlFlow Detected

    douglas125

      Hi.

      I'm trying to use a simple kernel to sum some values but I keep getting Error: Irreducible ControlFlow Detected with Stream SDK 2.1 on a Radeon 4890 Card. The kernel is attached. Any ideas?

      Thanks

       

       

      #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable #define MAX_GRIDSIZE 21 __kernel void CalcRowSignatures(__global uchar * ImgData, __global int * Region, __global int * DerivData, __global int * GridSize, __global float * Signature, __global int * width, __global float * FinalSignature) { int xfmx0 = Region[2]; int yfmy0 = Region[3]; int x = get_global_id(0); int gSize = GridSize[0]; int gSize2x=gSize*gSize*x; int valMedio = (gSize - 1) / 2; float RowSig[(MAX_GRIDSIZE+1)/2][MAX_GRIDSIZE]; for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { RowSig[i][j] = 0; } } for (int y = 0; y < yfmy0 - gSize; y++) { //Assinatura int corCentro = DerivData[x + valMedio + xfmx0 * (y + valMedio)]; for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { if (i != valMedio || j != valMedio) { int cor = DerivData[x + j + xfmx0 * (y + i)]; RowSig[i][j] += (float)cor * (float)corCentro; } } } } for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { Signature[gSize * i + j + gSize2x] = RowSig[i][j]; } } }

        • Irreducible ControlFlow Detected
          omkaranathan

          douglas,

          The issue has been informed to the developers and they are looking into it. Thanks for reporting

          • Irreducible ControlFlow Detected
            ryta1203

             

            Originally posted by: douglas125 Hi.

            I'm trying to use a simple kernel to sum some values but I keep getting Error: Irreducible ControlFlow Detected with Stream SDK 2.1 on a Radeon 4890 Card. The kernel is attached. Any ideas?

            Thanks

             

             

            This is an easy fix: change all your double for loops to single for loops. It might help to convert the 2D arrays to 1D arrays.

            Doing this allowed for compilation in SKA.

              • Irreducible ControlFlow Detected
                douglas125

                Fixed but I'm not very satisfied. The attached code works but it looks to me kinda redundant... I hope this is not what it should be like in the next versions.

                But thanks a lot for the help.

                 

                __kernel void CalcRowSignatures(__global uchar * ImgData, __global int * Region, __global int * DerivData, __global int * GridSize, __global float * Signature, __global int * width, __global float * FinalSignature) { int xfmx0 = Region[2]; int yfmy0 = Region[3]; int x = get_global_id(0); int gSize = GridSize[0]; int gSize2x=gSize*gSize*x; int valMedio = (gSize - 1) / 2; float RowSig[(MAX_GRIDSIZE+1)/2][MAX_GRIDSIZE]; for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { RowSig[i][j] = 0; } } //Assinatura for (int y = 0; y < yfmy0 - gSize; y++) { int corCentro = DerivData[x + valMedio + xfmx0 * (y + valMedio)]; int i = 0; int j = 0; for (int kk = 0; kk < (1+valMedio)*gSize; kk++) { if (i != valMedio || j != valMedio) { int cor = DerivData[x + j + xfmx0 * (y + i)]; RowSig[i][j] += (float)cor * (float)corCentro; } j++; if (j>=gSize) { j = 0; i++; } } } for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { Signature[gSize * i + j + gSize2x] = RowSig[i][j]; } } }

              • Irreducible ControlFlow Detected
                Raistmer
                Is there some workaround available? What about hotfix? Any ETA ?