be careful not to mix (early) return and barriers

A common practice (and who here has always a perfect [integer] global/local size ratio) is to create a padded global size and then within an opencl thread check for if the global_id() < some limit and return otherwise. What happens if there is a barrier later on in the code as shown below:

if global_id < some limit
    return
end

... do something ...

barrier(CLK_LOCAL_MEM_FENCE);


In this case the kernel may run or not. Specs say "All work-items in a work-group executing the kernel on a processor must execute this function before any are allowed to continue execution beyond the barrier. This function must be encountered by all work-items in a work-group executing the kernel."

http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/barrier.html

Either pad the arrays being acted on or do a conditional if global_id < some limit (run code) end everywhere

Remember the blocking in read/writes

A typical opencl example pseudo-code would be

some_func() -- assume that we are running on same context, valid kernel etc.
{
   allocate memory for a
   clEnqueueWriteBuffer(a)
   clEnqueueNDRangeKernel(using data in a)
   deallocate memory for a
}

would the above code always work if clEnqueueWriteBuffer(a) is a non-blocking write ?
not-necessarily due to the following

From http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueWriteBuffer.html

"blocking_write

Indicates if the write operations are blocking or nonblocking.

If blocking_write is CL_TRUE, the OpenCL implementation copies the data referred to by ptr and enqueues the write operation in the command-queue. The memory pointed to by ptr can be reused by the application after the clEnqueueWriteBuffer call returns.

If blocking_write is CL_FALSE, the OpenCL implementation will use ptr to perform a nonblocking write. As the write is non-blocking the implementation can return immediately. The memory pointed to by ptr cannot be reused by the application after the call returns. The eventargument returns an event object which can be used to query the execution status of the write command. When the write command has completed, the memory pointed to by ptr can then be reused by the application."

The assumption is the ptr (in this case data pointed by a) is available until the opencl command of write has finished. But if we are just queuing the commands and not blocking, a will be deallocated before its copied to the write buffer. The right thing to do in this case is to use a blocking write rather than a non-blocking write or use clFlush()/clFinish() before exiting the function.

Reads are also affected by blocking, so use blocking carefully. Especially be careful for that extra computation line in your function which may induce enough time for the non-blocking read to complete sometimes and not other times.


Ref:
http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueReadBuffer.html
http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clEnqueueWriteBuffer.html

bool data transfer and Opencl - Nah

Do not transfer bool data to an opencl kernel; instead convert it to int's or unsigned ints.

A big warning in https://www.khronos.org/registry/cl/api/2.0/cl.h

typedef cl_uint             cl_bool;                     /* WARNING!  Unlike cl_ types in cl_platform.h, cl_bool is not guaranteed to be the same size as the bool in kernels. */ 

You will get undefined results if you try transferring an array as bool.
http://devgurus.amd.com/thread/118829
http://stackoverflow.com/questions/4441280/does-opencl-support-boolean-variables