Hi, I’m trying to implement a simple Raytracing-Algorithm with OpenCl. This Algorithm shoud decide if a line is visible ore not.
My specs are
OS: Win8.1 with Microsoft Visual Studio 2010
Device name: Intel® HD Graphics 4600
Device version: OpenCL 1.2
I’m using the AMD APP SDK v2.9.1
The Build Error is
fcl build 1 succeeded.
error: Cannot select: 0x3c290a8: i64,ch = dynamic_stackalloc 0x3c2a0b0:1, 0x3c2a600, 0x3c29a38 [ORD=6] [ID=46]
0x3c2a600: i64 = bitcast 0x3c29928 [ID=45]
0x3c29928: v2i32 = IGILISD::MOVSWZ 0x3c2a2d0, 0x3c2a4f0, 0x3c2a710, 0x3c2a710 [ID=44]
0x3c2a2d0: i32 = shl 0x3c2af08, 0x3c29708 [ID=41]
0x3c2af08: i32 = IGILISD::LO_COPY 0x3c2f0d8 [ID=38]
0x3c2f0d8: i64 = bitcast 0x3c29e78 [ID=37]
0x3c29e78: v2i32 = IGILISD::MOVSWZ 0x3c299b0, 0x3c2a710, 0x3c2a710, 0x3c2a710 [ID=36]
0x3c299b0: i32 = IGILISD::LO_COPY 0x3c2a0b0 [ID=35]
0x3c2a0b0: i64,ch = load 0x3c2a798:1, 0x3c2aac8, 0x3c2abd8<LD4[%3], anyext from i32> [ID=34]
0x3c2aac8: i64 = add 0x3c2a1c0, 0x3c29ac0 [ORD=3] [ID=33]
0x3c2abd8: i64 = undef [ORD=4] [ID=5]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c29708: i32 = Constant<4> [ID=12]
0x3c2a4f0: i32 = or 0x3c2a468, 0x3c2a3e0 [ID=43]
0x3c2a468: i32 = shl 0x3c2a248, 0x3c29708 [ID=42]
0x3c2a248: i32 = IGILISD::HI_COPY 0x3c2f0d8 [ID=39]
0x3c2f0d8: i64 = bitcast 0x3c29e78 [ID=37]
0x3c29e78: v2i32 = IGILISD::MOVSWZ 0x3c299b0, 0x3c2a710, 0x3c2a710, 0x3c2a710 [ID=36]
0x3c299b0: i32 = IGILISD::LO_COPY 0x3c2a0b0 [ID=35]
0x3c2a0b0: i64,ch = load 0x3c2a798:1, 0x3c2aac8, 0x3c2abd8<LD4[%3], anyext from i32> [ID=34]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c29708: i32 = Constant<4> [ID=12]
0x3c2a3e0: i32 = srl 0x3c2af08, 0x3c2a358 [ID=40]
0x3c2af08: i32 = IGILISD::LO_COPY 0x3c2f0d8 [ID=38]
0x3c2f0d8: i64 = bitcast 0x3c29e78 [ID=37]
0x3c29e78: v2i32 = IGILISD::MOVSWZ 0x3c299b0, 0x3c2a710, 0x3c2a710, 0x3c2a710 [ID=36]
0x3c299b0: i32 = IGILISD::LO_COPY 0x3c2a0b0 [ID=35]
0x3c2a0b0: i64,ch = load 0x3c2a798:1, 0x3c2aac8, 0x3c2abd8<LD4[%3], anyext from i32> [ID=34]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a358: i32 = Constant<28> [ID=13]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c29a38: i64 = bitcast 0x3c2ad70 [ID=27]
0x3c2ad70: v2i32 = IGILISD::MOVSWZ 0x3c2a930, 0x3c2a710, 0x3c2a710, 0x3c2a710 [ID=21]
0x3c2a930: i32 = Constant<16> [ID=17]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
0x3c2a710: i32 = Constant<0> [ORD=1] [ID=3]
In function: Verdecktberechnung
error: midlevel compiler failed build.
Build failed!
The Code that cause the error is
bool AreaLineIntersection(float3 *Ecken, int arealen, float3 lineP1, float3 lineP2);
float3* PlanePoints(float3 *Ecken, int EckenLen, float3 *Erg);
float my_length(float3 v);
__kernel void Verdecktberechnung(__constant float3 *Ecken, __constant int *EckenLen, __constant int *offset, int len0, __constant float3 *line, __global int *Erg)
{
int id = get_global_id(0);
float3 lEcken[EckenLen[id]];
for(int i = 0; i < EckenLen[id]; i++)
lEcken[i] = Ecken[offset[id] + i];
float3 lLine[4];
for(int i = 0; i < 4; i++)
lLine[i] = line[i];
if(AreaLineIntersection(lEcken, EckenLen[id], lLine[0], lLine[1]) && AreaLineIntersection(lEcken, EckenLen[id], lLine[2], lLine[3]))
{
Erg[id] = 1;
}
else
{
Erg[id] = 0;
}
}
bool AreaLineIntersection(float3 *Ecken, int arealen, float3 lineP1, float3 lineP2)
{
if(arealen < 3)
return false;
float3 Erg[3];
PlanePoints(Ecken, arealen, Erg);
return true;
}
float3* PlanePoints(float3 *Ecken, int EckenLen, float3 *Erg)
{
if (EckenLen < 3)
{
Erg[0] = (float3)(-42,-42,-42);
return Erg;
}
if (EckenLen == 3)
{
Erg[0].x = Ecken[0].x;
Erg[0].y = Ecken[0].y;
Erg[0].z = Ecken[0].z;
Erg[1].x = Ecken[1].x;
Erg[1].y = Ecken[1].y;
Erg[1].z = Ecken[1].z;
Erg[2].x = Ecken[2].x;
Erg[2].y = Ecken[2].y;
Erg[2].z = Ecken[2].z;
return Erg;
}
if (EckenLen == 4)
{
float l1 = length(Ecken[1] - Ecken[0]);
float l2 = length(Ecken[2] - Ecken[1]);
float l3 = length(Ecken[3] - Ecken[2]);
float l4 = length(Ecken[0] - Ecken[3]);
if (l1 <= l2 && l1 <= l3 && l1 <= l4)
{
Erg[0].x = Ecken[0].x;
Erg[0].y = Ecken[0].y;
Erg[0].z = Ecken[0].z;
Erg[1].x = Ecken[2].x;
Erg[1].y = Ecken[2].y;
Erg[1].z = Ecken[2].z;
Erg[2].x = Ecken[3].x;
Erg[2].y = Ecken[3].y;
Erg[2].z = Ecken[3].z;
return Erg;
}
else if (l2 <= l1 && l2 <= l3 && l2 <= l4)
{
Erg[0].x = Ecken[0].x;
Erg[0].y = Ecken[0].y;
Erg[0].z = Ecken[0].z;
Erg[1].x = Ecken[1].x;
Erg[1].y = Ecken[1].y;
Erg[1].z = Ecken[1].z;
Erg[2].x = Ecken[3].x;
Erg[2].y = Ecken[3].y;
Erg[2].z = Ecken[3].z;
return Erg;
}
else
{
Erg[0].x = Ecken[0].x;
Erg[0].y = Ecken[0].y;
Erg[0].z = Ecken[0].z;
Erg[1].x = Ecken[1].x;
Erg[1].y = Ecken[1].y;
Erg[1].z = Ecken[1].z;
Erg[2].x = Ecken[2].x;
Erg[2].y = Ecken[2].y;
Erg[2].z = Ecken[2].z;
return Erg;
}
}
//float3 ptcopy[EckenLen];
//for(int i = 0; i < EckenLen; i++)
// ptcopy[i] = Ecken[i];
return &Erg[3];
}
The four lines
float l1 = length(Ecken[1] - Ecken[0]);
float l2 = length(Ecken[2] - Ecken[1]);
float l3 = length(Ecken[3] - Ecken[2]);
float l4 = length(Ecken[0] - Ecken[3]);
seems to cause the error.
With this lines comment out the Programm Build and Compile with out errors.
I have tryed to build it with
Kernel Builder for OpenCl API (Intel),
OpenCL Editor from OpenCLTemplate and
with a simple c++ host code which fails at Program.build().
It seems to be a bug in the AMD APP SDK, which shoud be fixed.
I try to Build it with the flag -cl-opt-disable,which solve the problem for some people. Not for me.
If i miss somthing please notice me and i will try to add them.
I very much apologize for my poor english, it is not my native-language.