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 ; } } }
douglas,
The issue has been informed to the developers and they are looking into it. Thanks for reporting
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.
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 ; } } }
And what was the correction?
I replaced the two loops in i and j with a loop in kk and manually incremented i and j as needed.
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.
Originally posted by: genaganna
For which devices, you are building this kernel.
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.
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