Em que minha GPU está esperando?

11

Estou escrevendo um programa OpenCL para uso com minha GPU AMD Radeon HD 7800 series. De acordo com o guia de programação OpenCL da AMD , essa geração de GPU possui duas filas de hardware que podem operar de forma assíncrona.

5.5.6 Fila de comandos

Para Ilhas do Sul e posteriores, os dispositivos suportam pelo menos duas filas de computação de hardware. Isso permite que um aplicativo aumente a taxa de transferência de pequenos despachos com duas filas de comandos para envio assíncrono e possivelmente execução. As filas de computação de hardware são selecionadas na seguinte ordem: primeira fila = filas de comandos OCL pares, segunda fila = filas OCL ímpares.

Para fazer isso, criei duas filas de comando OpenCL separadas para alimentar dados na GPU. Aproximadamente, o programa em execução no thread do host se parece com isso:

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

Com kNumQueues = 1, esse aplicativo funciona da maneira esperada: coleta todo o trabalho em uma única fila de comandos que é executada até a conclusão com a GPU estando bastante ocupada o tempo todo. Eu posso ver isso olhando a saída do criador de perfil CodeXL:

insira a descrição da imagem aqui

No entanto, quando eu defino kNumQueues = 2, espero que a mesma coisa aconteça, mas com o trabalho dividido igualmente em duas filas. De qualquer forma, espero que cada fila tenha as mesmas características individualmente que a fila: que comece a trabalhar sequencialmente até que tudo esteja pronto. No entanto, ao usar duas filas, posso ver que nem todo o trabalho está dividido nas duas filas de hardware:

insira a descrição da imagem aqui

No início do trabalho da GPU, as filas conseguem executar alguns kernels de forma assíncrona, embora pareça que nenhum deles ocupa totalmente as filas de hardware (a menos que meu entendimento esteja errado). Perto do final do trabalho da GPU, parece que as filas estão adicionando trabalho sequencialmente a apenas uma das filas de hardware, mas há até momentos em que nenhum kernel está em execução. O que da? Tenho algum mal-entendido fundamental de como o tempo de execução deve se comportar?

Eu tenho algumas teorias sobre por que isso está acontecendo:

  1. As clCreateBufferchamadas intercaladas estão forçando a GPU a alocar recursos de dispositivo de um conjunto de memórias compartilhadas de forma síncrona, que interrompe a execução de kernels individuais.

  2. A implementação OpenCL subjacente não mapeia filas lógicas para filas físicas e decide apenas onde colocar objetos em tempo de execução.

  3. Como estou usando objetos GL, a GPU precisa sincronizar o acesso à memória alocada especialmente durante as gravações.

Alguma dessas premissas é verdadeira? Alguém sabe o que pode estar causando a espera da GPU no cenário de duas filas? Todo e qualquer insight seria apreciado!

gpgpu opencl Mokosha
fonte
Não consigo encontrar onde diz que há duas filas de hardware no guia de programação. Você poderia postar uma cotação do documento? Mencione qual capítulo diz que há duas filas? O número de filas de hardware pode ser consultado em tempo de execução usando o OpenCL?
Andreas
Eu atualizei minha postagem. Diz a execução possível , mas se pode fazer alguns porque não pode fazê-los todos? Além disso, o tempo de execução do OpenCL não tem noção de fila de hardware, portanto não é algo que você possa consultar.
Mokosha 01/05/19

Respostas:

2

As filas de computação em geral não necessariamente significam que agora você pode fazer 2x despachos em paralelo. Uma única fila que sature totalmente as unidades de computação terá melhor rendimento. Várias filas são úteis se uma fila consome menos recursos (memória compartilhada ou registradores); as filas secundárias podem se sobrepor na mesma unidade de computação.

Para renderização em tempo real, esse é especialmente o caso de renderizações de sombra muito leves em computação / sombreadores, mas pesadas em hardware de função fixa, liberando assim o agendador de GPU para executar a fila secundária assíncrona.

Também encontrou isso nas notas de versão. Não sei se é o mesmo problema, mas pode ser que o CodeXL não seja ótimo. Eu esperaria que ele não tivesse a melhor instrumentação para a qual as expedições estão em voo.

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

Para um aplicativo que executa transferência de dados assíncrona simultânea e execução do kernel, a linha do tempo mostrada na visualização da sessão Rastreamento de Aplicativos não mostrará essas operações sobrepostas. Isso ocorre porque o driver e o hardware forçam essas operações a serem síncronas durante a criação de perfil. (333981)

Calvin
fonte