High Sierra, possible bug with compute kernels, intel HD 5000

Have each kernel write a value (say, 1) to an output buffer such that if all kernels run, the buffer is filled with 1s. Check it CPU-side. That should verify that all kernels have run. Did I understand the question correctly?

Yes, that's correct. The result is (current Xcode 9.0 and High Sierra on intel HD 5000):

1) for a short running kernel you get all 1's.

2) for a longer running kernel the output start with all 1's and the last part of the output-buffer all 0's.

So clearly not all kernels have run (I am speculating they got timed-out because the whole computation is about 15 seconds using a single dispatchThreadgroups on intel HD5000 [vs less than 3 seconds on my macpro Radeon R9 280X]), but there is no time-out error in the commandbuffer. Of course normally you don't know the output (that is why you're computing it in the first place), and the correctness of the compute result depends on the requirement that all kernels have run or that otherwise you get an error to inform the user that something went wrong.

So far I have only noticed this on intel HD 5000 but that might be because it is the slowest GPU I have tested on.

First thing that comes to mind is a memory-access error, but my kernel basically for each thread reads at 'thread_position_in_grid' and writes at 'thread_position_in_grid' and perfoms a for-loop of work. For for-loops less than 8192 iterations the total work is small enough that it works, but for larger amount of iterations it goes wrong. I do see on my mac-air visual artifacts in cocoa after this so that does indicate videocard-memory corruption. But so far I have not been able to figure out what is causing this. However, just splitting the work up in smaller work per thread using multiple dispatchThreadgroups and accumulating the result works fine.