I am writing an OpenCL program for use with my AMD Radeon HD 7800 series GPU. According to AMD's OpenCL programming guide, this generation of GPU has two hardware queues that can operate asynchronously.
5.5.6 Command Queue
For Southern Islands and later, devices support at least two hardware compute queues. That allows an application to increase the throughput of small dispatches with two command queues for asynchronous submission and possibly execution. The hardware compute queues are selected in the following order: first queue = even OCL command queues, second queue = odd OCL queues.
To do this, I've created two separate OpenCL command queues to feed data to the GPU. Roughly, the program running on the host thread looks something like this:
static const int kNumQueues = 2;
cl_command_queue default_queue;
cl_command_queue work_queue[kNumQueues];
static const int N = 256;
cl_mem gl_buffers[N];
cl_event finish_events[N];
clEnqueueAcquireGLObjects(default_queue, gl_buffers, N);
int queue_idx = 0;
for (int i = 0; i < N; ++i) {
cl_command_queue queue = work_queue[queue_idx];
cl_mem src = clCreateBuffer(CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, ...);
// Enqueue a few kernels
cl_mem tmp1 = clCreateBuffer(CL_READ_WRITE);
clEnqueueNDRangeKernel(kernel1, queue, src, tmp1);
clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp1);
cl_mem tmp2 = clCreateBuffer(CL_READ_WRITE);
clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp2);
clEnqueueNDRangeKernel(kernel3, queue, tmp2, gl_buffer[i], finish_events + i);
queue_idx = (queue_idx + 1) % kNumQueues;
}
clEnqueueReleaseGLObjects(default_queue, gl_buffers, N);
clWaitForEvents(N, finish_events);
With kNumQueues = 1
, this application pretty much works as intended: it collects all of the work into a single command queue that then runs to completion with the GPU being fairly busy the whole time. I'm able to see this by looking at the output of the CodeXL profiler:
However, when I set kNumQueues = 2
, I expect the same thing to happen but with the work evenly split across two queues. If anything, I expect each queue to have the same characteristics individually as the one queue: that it starts work sequentially until everything is done. However, when using two queues, I can see that not all of the work is split across the two hardware queues:
At the beginning of the GPU's work, the queues do manage to run some kernels asynchronously, although it seems like neither ever fully occupies the hardware queues (unless my understanding is mistaken). Near the end of the GPU work, it seems like the queues are adding work sequentially to only one of the hardware queues, but there are even times that no kernels are running. What gives? Do I have some fundamental misunderstanding of how the runtime is supposed to behave?
I have a few theories as to why this is happening:
The interspersed
clCreateBuffer
calls are forcing the GPU to allocate device resources from a shared memory pool synchronously which stalls the execution of individual kernels.The underlying OpenCL implementation does not map logical queues to physical queues, and only decides where to place objects at runtime.
Because I'm using GL objects, the GPU needs to synchronize access to specially allocated memory during writes.
Are any of these assumptions true? Does anyone know what could be causing the GPU to wait in the two-queue scenario? Any and all insight would be appreciated!