What is my GPU waiting on?

Question

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:

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!

Answer 1

Compute queues in general don't necessarily mean you can now do 2x dispatches in parallel. A single queue that fully saturates the compute units will have better throughput. Multiple queues are useful if one queue consumes less resources (shared memory or registers), then secondary queues can then overlap on the same compute unit.

For real-time rendering this is especially the case with things like shadow rendering that are very light on compute/shaders but heavy on fixed function hardware, thus freeing the GPU scheduler to run the secondary queue async.

Also found this in the release notes. Don't know if it's the same issue, but could just be that CodeXL is not great. I would expect that it might have not have the best instrumentation for which dispatches are in flight.

https://developer.amd.com/wordpress/media/2013/02/AMD_CodeXL_Release_Notes.pdf

For an application that performs simultaneous asynchronous data transfer and kernel execution, the timeline shown in the Application Trace session view will not show these operations being overlapped. This is because the driver and hardware force these operations to be synchronous while profiling. (333981)