Sample kernel showing multiple driver bugs (printf-related?)

[Bruce_M_]

The attached code produces some very weird results, that I think must be OpenCL compiler bugs. I tried to make it smaller but it seems sensitive to small changes (it's still a pretty minimal example though).

The output is

[plain]Device name:         Intel(R) Core(TM) i7-2600 CPU @ 3.40GHz
Driver version: 1.2.0.76921
addr = 0 lid = 0 bias = 0
addr = 1 lid = 1 bias = 0
addr = 2 lid = 2 bias = 0
addr = 3 lid = 3 bias = 0
addr = 4 lid = 4 bias = 0
addr = 5 lid = 5 bias = 0
addr = 6 lid = 6 bias = 0
addr = 7 lid = 7 bias = 0
addr = 8 lid = 8 bias = 0
addr = 9 lid = 9 bias = 0
addr = 10 lid = 10 bias = 0
addr = 11 lid = 11 bias = 0
addr = 12 lid = 12 bias = 0
addr = 13 lid = 13 bias = 0
addr = 14 lid = 14 bias = 0
addr = 15 lid = 15 bias = 0
addr = 16 lid = 0 bias = 16
addr = 16 lid = 1 bias = 16
addr = 16 lid = 2 bias = 16
addr = 16 lid = 3 bias = 16
addr = 16 lid = 4 bias = 16
addr = 16 lid = 5 bias = 16
addr = 16 lid = 6 bias = 16
addr = 16 lid = 7 bias = 16
addr = 16 lid = 8 bias = 16
addr = 16 lid = 9 bias = 16
addr = 16 lid = 10 bias = 16
addr = 16 lid = 11 bias = 16
addr = 16 lid = 12 bias = 16
addr = 16 lid = 13 bias = 16
addr = 16 lid = 14 bias = 16
addr = 16 lid = 15 bias = 16
00 00 00 00 c4 7f 00 00 40 66 f0 b6 c4 7f 00 00 00 00 00 00 01 00 00 00 4d 0f 4d bc c4 7f 00 00[/plain]

There are several oddities here:

1. For the second half, addr is reported as 16, even though it is computed as lid + bias, which should range from 16 to 31.
2. The hex dump line (final line) only appears once, even though upsweep4 is called twice.
3. The hex dump shows random garbage (changes from run to run), even though the two calls to upsweepMulti should have zeroed out the entire array.

As a bonus bug, uncommenting the #define OPT_DISABLE in the host code (to compile with -cl-opt-disable) changes the results such that get_local_id(0) appears to be returning 0 everywhere.

I'm using 64-bit Ubuntu 13.10. Hopefully the bug will reproduce on other systems, given that it seems like a memory corruption bug of some sort.

[Raghupathi_M_Intel]

We will take a look at the issues and get back to you.

Thanks,
Raghu