OpenCL race condition with clSetKernelArg

user_11757687 0
Viewed 460

From the Khronos website on the thread safety of clSetKernelArg:

All OpenCL API calls are thread-safe except clSetKernelArg, which is safe to call from any host thread, and is safe to call re-entrantly so long as concurrent calls operate on different cl_kernel objects. However, the behavior of the cl_kernel object is undefined if clSetKernelArg is called from multiple host threads on the same cl_kernel object at the same time.

My question is, is there a way to make this behavior defined where kernels can read and write from a single kernel object from multiple threads?

I considered that std::atomic on the object being modified by the kernels would prevent this undefined behavior, but from what I have tried, it results in the kernel's output producing the wrong values. Is there a better way to implement this/ a known technique for dealing with a case?

It might be useful in a case where the allocated object's size is so large that recreating a new object for every kernel execution costs too much memory, and a shared/overridable object would be preferred.

opencl kernel race-condition multithreading c++
1 Answers

where kernels can read and write from a single kernel object from multiple threads?

By "kernels" you mean pieces of code executing on GPU, and by "single kernel object" you mean cl_kernel in host code ? Kernels on GPU don't ever see the cl_kernel struct that exists on the host side. I assume you're talking about using buffer object (cl_mem) arguments by kernels.

You can think of cl_kernel as:

struct {
  size_t num_args;
  void* args[];
} _cl_kernel;
typedef struct _cl_kernel * cl_kernel;

If you call clSetKernelArg(), it simply sets something in that struct. If you call clEnqueueNDRangeKernel(), it takes a snapshot of the cl_kernel struct (of the arguments), and appends it to some internal device queue. By "snapshot" i don't mean it creates a hidden snapshot of a actual cl_mem buffer contents; it simply copies a reference to the cl_mem arguments. Since it's a reference, it doesn't really matter whether you use a single cl_kernel object from multiple threads, or you call clCreateKernel multiple times with same name, and then use those cl_kernel in each thread; it's just a matter of convenience, end result is the same.

If you have a single in-order command queue, your kernels will execute deterministically, in enqueue order. If you have multiple command queues (in- or out-of-order, doesn't matter), there isn't any implicit ordering between the queues, so if you enqueue the same kernel into all of the queues, they'll execute in random order. You can force an explicit order by using events. IOW, you do:

cl_event event1, event2;
cl_kernel K;
...
clEnqueueNDRangeKernel(queue_1, K, ... , &event1);
clEnqueueNDRangeKernel(queue_2, K, ... , 1, &event1, &event2);

etc. This will force a kernel execution to wait on previous one, even if they're on different queues. But you'll only have one kernel using the buffer at a time.

If you want to use the same buffer simultaneously by multiple running kernels, then it depends or usage pattern of that buffer. If you're doing only reads, you can safely use the buffer from any number of kernels simultaneously. For write use, if you know you'll only write into a part of the buffer, you could try using sub-buffers (clCreateSubBuffer). Otherwise you're probably out of luck (well maybe you could try atomic ops, but it could make an algorithm unusably slow).

user_57677250
Related