¿Qué está esperando mi GPU?

11

Estoy escribiendo un programa OpenCL para usar con mi GPU AMD Radeon HD 7800 series. Según la guía de programación OpenCL de AMD , esta generación de GPU tiene dos colas de hardware que pueden funcionar de forma asíncrona.

5.5.6 Cola de comandos

Para las Islas del Sur y posteriores, los dispositivos admiten al menos dos colas de cómputo de hardware. Eso permite que una aplicación aumente el rendimiento de pequeños despachos con dos colas de comandos para el envío asíncrono y posiblemente la ejecución. Las colas de cálculo de hardware se seleccionan en el siguiente orden: primera cola = colas de comandos OCL pares, segunda cola = colas OCL impares.

Para hacer esto, he creado dos colas de comandos OpenCL separadas para alimentar datos a la GPU. Aproximadamente, el programa que se ejecuta en el subproceso host se parece a esto:

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);

Con kNumQueues = 1, esta aplicación funciona más o menos según lo previsto: recopila todo el trabajo en una sola cola de comandos que luego se ejecuta hasta su finalización con la GPU bastante ocupada todo el tiempo. Puedo ver esto mirando la salida del generador de perfiles CodeXL:

ingrese la descripción de la imagen aquí

Sin embargo, cuando configuro kNumQueues = 2, espero que suceda lo mismo pero con el trabajo dividido en dos colas. En todo caso, espero que cada cola tenga las mismas características individualmente que la cola única: que comience a funcionar secuencialmente hasta que todo esté hecho. Sin embargo, cuando uso dos colas, puedo ver que no todo el trabajo se divide en las dos colas de hardware:

ingrese la descripción de la imagen aquí

Al comienzo del trabajo de la GPU, las colas logran ejecutar algunos núcleos de forma asincrónica, aunque parece que ninguno de ellos ocupa completamente las colas de hardware (a menos que mi comprensión esté equivocada). Cerca del final del trabajo de GPU, parece que las colas están agregando trabajo secuencialmente a solo una de las colas de hardware, pero incluso hay veces que no se están ejecutando núcleos. ¿Lo que da? ¿Tengo algún malentendido fundamental sobre cómo se supone que debe comportarse el tiempo de ejecución?

Tengo algunas teorías sobre por qué sucede esto:

  1. Las clCreateBufferllamadas intercaladas obligan a la GPU a asignar recursos del dispositivo desde un conjunto de memoria compartida de forma síncrona que detiene la ejecución de núcleos individuales.

  2. La implementación subyacente de OpenCL no asigna las colas lógicas a las colas físicas, y solo decide dónde colocar los objetos en tiempo de ejecución.

  3. Como estoy usando objetos GL, la GPU necesita sincronizar el acceso a la memoria especialmente asignada durante las escrituras.

¿Alguno de estos supuestos es cierto? ¿Alguien sabe qué podría estar causando que la GPU espere en el escenario de dos colas? Cualquier apreciación será apreciada!

gpgpu opencl Mokosha
fuente
No puedo encontrar dónde dice que hay dos colas de hardware en la guía de programación. ¿Podría publicar una cita del documento? ¿Menciona qué capítulo que dice que hay dos colas? ¿Se puede consultar el número de colas de hardware en tiempo de ejecución con OpenCL?
Andreas
He actualizado mi publicación. Dice posible ejecución, pero si puede hacer algunas, ¿por qué no puede hacerlas todas? Además, el tiempo de ejecución de OpenCL no tiene noción de cola de hardware, por lo que no es algo que pueda consultar.
Mokosha

Respuestas:

2

Calcular colas en general no necesariamente significa que ahora puede hacer 2x despachos en paralelo. Una sola cola que satura completamente las unidades de cómputo tendrá un mejor rendimiento. Las colas múltiples son útiles si una cola consume menos recursos (memoria compartida o registros), luego las colas secundarias pueden superponerse en la misma unidad de cómputo.

Para la representación en tiempo real, este es especialmente el caso con cosas como la representación de sombras que son muy ligeras en cómputo / sombreadores pero pesadas en hardware de función fija, liberando así el programador de GPU para ejecutar la cola secundaria asíncrona.

También encontré esto en las notas de la versión. No sé si es el mismo problema, pero podría ser que CodeXL no es genial. Esperaría que no tuviera la mejor instrumentación para los despachos en vuelo.

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

Para una aplicación que realiza transferencia de datos asincrónica simultánea y ejecución de kernel, la línea de tiempo que se muestra en la vista de sesión de Application Trace no mostrará que estas operaciones se superponen. Esto se debe a que el controlador y el hardware obligan a estas operaciones a ser síncronas durante la creación de perfiles. (333981)

Calvin
fuente