this part of my code
if (get_global_id(0)==TRACE_TID) printf((__constant char *)"div2.6: q.d4=%x, carry=%x, nn.d3=%x\n", q.d4, carry, nn.d3);
q.d4 = q.d4 - nn.d3 + carry;
if (get_global_id(0)==TRACE_TID) printf((__constant char *)"div2.7: q.d4=%x, carry=%x, nn.d3=%x\n", q.d4, carry, nn.d3);
prints the following:
div2.6: q.d4=51ed, carry=0, nn.d3=51ed
div2.7: q.d4=394d8646, carry=0, nn.d3=51ed
So, according to Catalyst 13.9 on HD5770,
0x51ed - 0x51ed + 0 = 0x394d8646.
In fact each time I run it, the result is something different. All variables/struct members are of type uint. 13.10 beta does not correct the problem. When I redefine the variables' type to uint2 or uint4, the code works correctly. It also works fine when running it on the CPU device.
Of course, when just putting that into a small test program, it always works. I can provide all source code and/or binaries if needed. But first I wanted to check if that is a known problem?