OpenCL parse error, only on GPU device

Question

When compiling the OpenCL kernel below to my GPU (HD Graphics 5000) I get "Parse error." from PROGRAM_BUILD_LOG.

constant int cols = 946;

kernel void run(global const uchar4 *curr) {
    int row = get_global_id(0);
    int col = get_global_id(1);

    int cell = row * cols + col;

    if (col > 0 && col < (cols - 1)) {
        if (curr[ cell ].x == 243) {
            // something
        }
    }
}

However when compiling to my CPU it works fine. Also, by changing very little in the code example it will compile just fine. Here are 3 examples there all works.

Example 1:

constant int cols = 946;

kernel void run(global const uchar4 *curr) {
    int row = get_global_id(0);
    int col = get_global_id(1);

    if (col > 0 && col < (cols - 1)) {
        if (curr[ row * cols + col ].x == 243) {
            // something
        }
    }
}

Example 2:

kernel void run(global const uchar4 *curr) {
    int row = get_global_id(0);
    int col = get_global_id(1);
    int cols = 946;

    int cell = row * cols + col;

    if (col > 0 && col < (cols - 1)) {
        if (curr[ cell ].x == 243) {
            // something
        }
    }
}

Example 3:

constant int cols = 946;

kernel void run(global const uchar4 *curr) {
    int row = get_global_id(0);
    int col = get_global_id(1);

    int cell = row * cols + col;

    if (curr[ cell ].x == 243) {
        // something
    }
}

Answer 1

I've just tried to compile this code for a HD 4000 on OS X and receive the same error. Given the nature of the build log and the fact that the same code builds successfully on other devices, this is clearly a bug with Apple's OpenCL implementation. In my experience Apple's OpenCL implementation exhibits a particularly large number of bugs, often involving compilation failures with useless error messages. The HD graphics devices seem to be responsible for a significant number of these (this is the 3rd bug for HD graphics on OS X posted to stack overflow in the last two weeks), perhaps because the implementation for them is still relatively immature.

I recommend you raise a bug via the Apple Bug Reporting System.