OpenCL crash when Max work group size is used

Question

So I have a problem.

I am writing an application which uses OpenCL and whenever I use the max work group size or anything above half of the max work group size, I get a crash (a black screen).

Does anyone know what the reason might be?

I would love to use to entire work group size, instead of just half.

Thanks in advance

Answer 1

The most likely reason is that you have a global array of a size that, at some point, is not evenly divisible by the workgroup size.

For example a global array of size 1920. Workgroup size 32/64/128 works, then you get 60/30/15 workgroups. But if you choose workgroup size 256 (or larger), you get 7.5 workgroups. What then happens is that 8 workgropus are executed, and the last workgroup covers threads 1792-2047. Yet, the global array ends at tread 1919, so access to array elements 1920-2047 reads from / writes to Nirwana and this can crash the entire program.

There are 2 possible solutions:

1. Only choose a workgroup size small enough that all global array sizes are divisible by it. Note that workgroup size must be 32 or a multiple of 32.
2. Alternatively, at the very begin of the Kernel code, use a so-called "guard clause": if(get_global_id(0)>=array_size) return;. This makes the thread end if the there is no corresponding memory address to the thread ID, preventing crashing. With this, you can also use any array size that is not divisible by the workgroup size.