Archives Discussions

[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 = 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 += (float)cor * (float)corCentro; } } } } for (int i = 0; i <= valMedio; i++) { for (int j = 0; j < gSize; j++) { Signature[gSize * i + j + gSize2x] = RowSig; } } }

[omkaranathan]

douglas,

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

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

[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 = 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 += (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; } } }

[Lev]

And what was the correction?

[douglas125]

I replaced the two loops in i and j with a loop in kk and manually incremented i and j as needed.

[genaganna]

Originally posted by: Raistmer After upgrading to SDK 2.2 and Cat 10.7 I started to get this error: Error: Building Program (clBuildProgram):main kernels: -11 Error:E010:Irreducible ControlFlow Detected I didn't recive this error before. Moreover, when CL file compiled via clc from console - I get no errors! This error appears only when kernels are compiled from application. What solution? EDIT: No errors if attached kernel is disabled What wrong with this kernel and why clc doesn't show any errors when file compiled from command line? 😕

For which devices,  you are building this kernel.

[Raistmer]

Originally posted by: genaganna

For which devices,  you are building this kernel.

For HD4870.

[genaganna]

Originally posted by: genaganna

Originally posted by: Raistmer After upgrading to SDK 2.2 and Cat 10.7 I started to get this error: Error: Building Program (clBuildProgram):main kernels: -11 Error:E010:Irreducible ControlFlow Detected I didn't recive this error before. Moreover, when CL file compiled via clc from console - I get no errors! This error appears only when kernels are compiled from application. What solution? EDIT: No errors if attached kernel is disabled What wrong with this kernel and why clc doesn't show any errors when file compiled from command line? 😕

For which devices,  you are building this kernel.

We are able to reproduce this issue.  reported to developers.

[Raistmer]

Is there some workaround available? What about hotfix? Any ETA ?

[ryta1203]

Originally posted by: Raistmer Is there some workaround available? What about hotfix? Any ETA ?

Have you tried eliminating all those unnecessary nested ifs?

1. Use the select()

2. Just use one if, combine the conditional

3. Use software-based predication