This kernel below crashes the compiler.
Not only only that but during the process it causes some sort of internal compiler debug text to be printed on my own program console window (!!) hundreds of lines of myserious internal debug code...not sure how...
The obvious cause of the error is the lines "BSV0 &= InsideMask" and "BSV1 &= InsideMask;" (which I know are incorrect)...the compiler doesn't crash with these removed.
The program crashes after printing this text, so I can't copy all of it, but it's along the lines of:
0x6758540: v2f32,ch,glue = CopyFromReg 0x6757300, 0x675adc0, 0x6757308:1 [ ORD=1877] [ID=290]
0x675aba0: i32 = TargetConstant<1> [ORD=1761] [ID=38]
Windows 7 64bit, ATI 5870, SDK 2.6, Catalyst 12.2 preview
|Here's the code|
#define InfinityCorrectfloat3(val) as_float3(as_uint3(val) - (((as_uint3(val) & 0x7F800000) + 1)>>31))
inline float3 safe_recip_float3(const float3 Val)
__kernel void Test(__constant float* TraceParamsD, __global const int* TraceParamsI,__global const uint *ATree,__write_only image2d_t ScreenBuffer)
PixelPos.x = TraceParamsI[32+get_global_id(0)];
if(NodesToCheck == 0)
for(TriID = ATree[NodeOffset+11]-1; TriID >= 0; TriID--)
VertX = mad(OVertZ,sinY,OVertX*cosY);
val = (VertX.s3 * VertY.s1) - (VertX.s3 * VertY.s2) - (VertX.s1 * VertY.s3) + (VertX.s2 * VertY.s3) + val2;
Thanks for reporting this issue. I've been able to reproduce it and have reported it to the correct person to investigate. The problem seems to stem from operation.
BSV0 &= InsideMask; <--
This is making it past our semantic checker and should not as BSV0 is a floating point operation and thus a binary operand is not legal. Please switch it to BSV0 = as_float(as_int(BSV0) & InsideMask); as a work-around until we release a fix for the crash.