What is my GPU waiting on?

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
clEnqueueNDRangeKernel(kernel1, queue, src, tmp1);

clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp1);

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:

1. 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.

2. The underlying OpenCL implementation does not map logical queues to physical queues, and only decides where to place objects at runtime.

3. 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!

• I can't find where it says there are two hardware queues in the programming guide. Could you post a quote from the document? Mention which chapter that says there are two queues? Is number of hardware queues queryable in runtime using OpenCL? May 1, 2016 at 12:06
• I've updated my post. It does say possible execution, but if it can do a few why can't it do them all? Also the OpenCL runtime has no notion of hardware queue, so it's not something that you can query. May 1, 2016 at 16:05