Get best of the week

Computação GPU - Por que, quando e como. Mais alguns testes

Todos sabem há muito tempo que, nas placas de vídeo, você não pode apenas brincar com brinquedos, mas também executar coisas que não estão relacionadas a jogos, por exemplo, treinar uma rede neural, lembrar-se de criptomoeda ou realizar cálculos científicos. Como aconteceu, você pode lê-lo aqui , mas eu queria falar sobre o motivo pelo qual a GPU pode ser interessante para o programador médio (não relacionado ao GameDev) como abordar o desenvolvimento na GPU sem gastar muito tempo, decidir se olhe nessa direção e " descubra com os dedos" qual o lucro que você pode obter. 



O artigo foi escrito com base na minha apresentação no HighLoad ++. Ele discute principalmente as tecnologias oferecidas pela NVIDIA. Não tenho o propósito de anunciar produtos, apenas os cito como exemplo e, com certeza, algo semelhante pode ser encontrado em fabricantes concorrentes.

Por que contar com a GPU?


Dois processadores podem ser comparados de acordo com critérios diferentes, provavelmente os mais populares são a frequência e o número de núcleos, o tamanho dos caches etc., mas no final, estamos interessados ​​em quantas operações um processador pode executar por unidade de tempo, que tipo de operação é essa, mas uma pergunta separada Uma métrica comum é o número de operações de ponto flutuante por segundo - flops. E quando queremos comparar quente com macio e, no nosso caso, GPU com CPU, essa métrica é útil.

O gráfico abaixo mostra o crescimento desses mesmos fracassos ao longo do tempo para processadores e placas de vídeo.


(Os dados são coletados de fontes abertas, não há dados para 2019-20 anos, porque nem tudo é tão bonito lá, mas as GPUs ainda vencem)

Bem, é tentador, não é? Mudamos todos os cálculos da CPU para a GPU e obtemos oito vezes o melhor desempenho!

Mas, claro, nem tudo é tão simples. Você não pode simplesmente levar e transferir tudo para a GPU, por que conversaremos mais.

Arquitetura da GPU e sua comparação com a CPU


Trago para muitos uma imagem familiar com a arquitetura da CPU e os elementos básicos:


Núcleo da CPU

O que há de tão especial? Um núcleo e um monte de blocos auxiliares.

Agora vamos ver a arquitetura da GPU:


Núcleo da GPU

Uma placa de vídeo possui muitos núcleos de processamento, geralmente vários milhares, mas são combinados em blocos; para placas de vídeo NVIDIA, geralmente 32 cada, e possuem elementos comuns, incluindo e registros. A arquitetura do núcleo da GPU e dos elementos lógicos é muito mais simples do que na CPU, ou seja, não há pré-buscadores, preditores de brunch e muito mais.

Bem, esses são os pontos principais da diferença na arquitetura da CPU e da GPU e, de fato, impõem restrições ou, inversamente, abrem as possibilidades do que podemos ler efetivamente na GPU.

Eu não mencionei mais um ponto importante, geralmente, a placa de vídeo e o processador não "remexem" entre si e gravam dados na placa de vídeo e lê o resultado de volta - essas são operações separadas e podem vir a ser um "gargalo" em seu sistema, um gráfico do tempo de bombeamento versus tamanho dados são fornecidos posteriormente neste artigo.

Limitações e recursos da GPU


Quais limitações essa arquitetura impõe aos algoritmos executáveis:

  • Se estivermos calculando em uma GPU, não podemos selecionar apenas um núcleo, um bloco inteiro de núcleos será alocado (32 para NVIDIA).
  • Todos os núcleos executam as mesmas instruções, mas com dados diferentes (falaremos sobre isso mais adiante), esses cálculos são chamados de instrução única - dados múltiplos ou SIMD (embora a NVIDIA introduza seu refinamento). 
  • Devido ao conjunto relativamente simples de blocos lógicos e registros gerais, a GPU realmente não gosta de ramificação e, de fato, a lógica complexa nos algoritmos.

Que oportunidades ele abre:

  • Na verdade, a aceleração desses mesmos cálculos SIMD. O exemplo mais simples é a adição elementar de matrizes, e vamos analisá-lo.

Redução de algoritmos clássicos para representação SIMD


Transformação


Temos duas matrizes, A e B, e queremos adicionar um elemento da matriz B a cada elemento da matriz A. Abaixo está um exemplo em C, embora eu espero que fique claro para aqueles que não falam esse idioma:

void func(float *A, float *B, size)
{ 
   for (int i = 0; i < size; i++) 
   { 
       A[i] += B[i]
   } 
}

Loopback clássico de elementos em um loop e tempo de execução linear.

Agora vamos ver como esse código ficará para a GPU:

void func(float *A, float *B, size) 
{ 
   int i = threadIdx.x; 
   if (i < size) 
      A[i] += B[i] 
}

E aqui já é interessante, a variável threadIdx apareceu, o que não parecemos declarar em lugar algum. Sim, seu sistema nos fornece. Imagine que no exemplo anterior a matriz consiste em três elementos e você deseja executá-la em três threads paralelos. Para fazer isso, você precisaria adicionar outro parâmetro - o número do índice ou fluxo. É isso que a placa de vídeo faz por nós, embora ela passe o índice como uma variável estática e possa trabalhar com várias dimensões ao mesmo tempo - x, y, z.

Outra nuance: se você iniciar um grande número de fluxos paralelos de uma só vez, os fluxos deverão ser divididos em blocos (um recurso arquitetônico das placas de vídeo). O tamanho máximo do bloco depende da placa de vídeo, e o índice do elemento para o qual realizamos cálculos precisará ser obtido da seguinte maneira:

int i = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx –  , blockDim –  , threadIdx –

Como resultado, o que temos: muitos threads de execução paralela que executam o mesmo código, mas com índices diferentes e, consequentemente, dados, ou seja, o mesmo SIMD.

Este é o exemplo mais simples, mas se você quiser trabalhar com a GPU, precisará levar sua tarefa para o mesmo formulário. Infelizmente, isso nem sempre é possível e, em alguns casos, pode se tornar objeto de uma dissertação de doutorado, mas, no entanto, algoritmos clássicos ainda podem ser trazidos para esse formulário.

Agregação


Vamos agora ver como será a agregação convertida para a representação SIMD:
 

Temos uma matriz de n elementos. No primeiro estágio, iniciamos n / 2 threads e cada thread adiciona dois elementos, ou seja, em uma iteração, adicionamos metade dos elementos na matriz. E então, no loop, repetimos a mesma coisa para a matriz recém-criada, até agregar os dois últimos elementos. Como você pode ver, quanto menor o tamanho da matriz, menos threads paralelos podemos iniciar, ou seja, em uma GPU, faz sentido agregar matrizes de tamanho suficientemente grande. Esse algoritmo pode ser usado para calcular a soma de elementos (a propósito, não se esqueça do possível estouro do tipo de dados com o qual você está trabalhando), procure o máximo, o mínimo ou apenas a pesquisa.

Ordenação


Mas a classificação já parece muito mais complicada.

Os dois algoritmos de classificação mais populares na GPU são:

  • Bitonic-sort
  • Radix-sort

Porém, o radix-sort ainda é usado com mais frequência e a implementação pronta para produção pode ser encontrada em algumas bibliotecas. Não analisarei detalhadamente como esses algoritmos funcionam; os interessados ​​podem encontrar uma descrição do tipo de raiz em https://www.codeproject.com/Articles/543451/Parallel-Radix-Sort-on-the-GPU-using-Cplusplus- AMP e https://stackoverflow.com/a/26229897

Mas a ideia é que mesmo um algoritmo não linear como a classificação possa ser reduzido a uma visualização SIMD.

E agora, antes de analisar os números reais que podem ser obtidos na GPU, vamos descobrir como programar para esse milagre da tecnologia?

Onde começar


As duas tecnologias mais comuns que podem ser usadas para programação na GPU:

  • Opencl
  • Cuda

O OpenCL é um padrão suportado pela maioria dos fabricantes de placas de vídeo, incluindo e em dispositivos móveis, também o código escrito em OpenCL pode ser executado na CPU.

Você pode usar o OpenCL a partir de C / C ++, existem pastas para outros idiomas.

Para o OpenCL, gostei mais do livro OpenCL in Action . Também descreve diferentes algoritmos na GPU, incluindo Classificação Bitonic e Radix.

CUDA é a tecnologia proprietária e o SDK da NVIDIA. Você pode escrever em C / C ++ ou usar ligações para outros idiomas.

Comparar OpenCL e CUDA não está correto, porque um é o padrão, o outro é o SDK inteiro. No entanto, muitas pessoas escolhem a CUDA para o desenvolvimento de placas de vídeo, apesar de a tecnologia ser proprietária, embora gratuita e funcionando apenas em placas NVIDIA. Há várias razões para isso:

  • API
  • , GPU, (host)
  • , ..

As peculiaridades incluem o fato de que o CUDA vem com seu próprio compilador, que também pode compilar o código C / C ++ padrão.

O livro CUDA mais abrangente que me deparei foi a Programação CUDA C profissional , embora já esteja um pouco desatualizada, mas discute muitas nuances técnicas de programação para placas NVIDIA.

Mas e se eu não quiser passar alguns meses lendo esses livros, escrevendo meu próprio programa para uma placa de vídeo, testando e depurando e depois descobrindo que isso não é para mim? 

Como eu disse, há um grande número de bibliotecas que escondem a complexidade do desenvolvimento sob a GPU: XGBoost, cuBLAS, TensorFlow, PyTorch e outras, consideraremos a biblioteca de empuxo, uma vez que é menos especializada do que as outras bibliotecas acima, mas ao mesmo tempo implementa algoritmos básicos, por exemplo, classificação, pesquisa, agregação e com alta probabilidade de aplicabilidade em suas tarefas.

Thrust é uma biblioteca C ++ que visa "substituir" algoritmos STL padrão por algoritmos baseados em GPU. Por exemplo, classificar uma matriz de números usando esta biblioteca em uma placa de vídeo seria assim:

thrust::host_vector<DataType> h_vec(size); //    
std::generate(h_vec.begin(), h_vec.end(), rand); //   
thrust::device_vector<DataType> d_vec = h_vec; //         
thrust::sort(d_vec.begin(), d_vec.end()); //    
thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); //   ,

(não esqueça que o exemplo deve ser compilado por um compilador da NVIDIA)

Como você pode ver, o thrust :: sort é muito semelhante a um algoritmo semelhante do STL. Esta biblioteca esconde muitas dificuldades, em particular o desenvolvimento de um subprograma (mais precisamente, o kernel), que será executado na placa de vídeo, mas, ao mesmo tempo, priva a flexibilidade. Por exemplo, se quisermos classificar vários gigabytes de dados, seria lógico enviar um dado para o cartão para iniciar a classificação e, enquanto a classificação estiver em andamento, envie mais dados para o cartão. Essa abordagem é chamada ocultação de latência e permite um uso mais eficiente dos recursos de mapas do servidor, mas, infelizmente, quando usamos bibliotecas de alto nível, essas oportunidades permanecem ocultas. Mas para prototipar e medir o desempenho, eles são iguais, especialmente com o empuxo, que você pode medir qual sobrecarga a transferência de dados fornece.

Eu escrevi uma pequena referência usando esta biblioteca, que executa vários algoritmos populares com diferentes quantidades de dados na GPU, vamos ver quais são os resultados.

Resultados do algoritmo GPU


Para testar a GPU, eu instalei uma instância na AWS com uma placa de vídeo Tesla k80, essa não é a placa de servidor mais poderosa até o momento (a Tesla v100 mais poderosa), mas a mais acessível e a bordo:

  • 4992 CUDA Kernels
  • 24 GB de memória
  • 480 Gb / s - largura de banda de memória 

E para testes na CPU, tomei uma instância com um processador Intel Xeon CPU E5-2686 v4 a 2.30GHz

Transformação



Tempo de execução da transformação na GPU e CPU em ms

Como você pode ver, a transformação usual dos elementos da matriz é aproximadamente a mesma no tempo, tanto na GPU quanto na CPU. E porque? Como a sobrecarga para enviar dados para o cartão e vice-versa consome todo o aumento de desempenho (falaremos sobre sobrecarga separadamente), e há relativamente poucos cálculos no cartão. Além disso, não esqueça que os processadores também suportam instruções SIMD e os compiladores, em casos simples, podem usá-los efetivamente. 

Vamos agora ver com que eficiência a agregação é feita na GPU.

Agregação



Tempo de execução da agregação na GPU e CPU em ms

No exemplo de agregação, já observamos um aumento significativo no desempenho com um aumento no volume de dados. Também vale a pena prestar atenção ao fato de bombearmos uma grande quantidade de dados para a memória do cartão e apenas um valor agregado é recuperado, ou seja, A sobrecarga para transferir dados do cartão para a RAM é mínima.

Vamos para o exemplo mais interessante - a classificação.

Ordenação



Classificação do tempo na GPU e CPU em ms

Apesar de enviarmos toda a matriz de dados para a placa de vídeo e vice-versa, a classificação na GPU de 800 MB de dados é aproximadamente 25 vezes mais rápida que no processador.

Sobrecarga de transferência de dados


Como pode ser visto no exemplo de transformação, nem sempre é óbvio se a GPU será eficaz, mesmo nas tarefas que são paralelas. O motivo disso é uma sobrecarga para transferir dados da RAM do computador para a memória da placa de vídeo (em consoles de jogos, a propósito, a memória é compartilhada entre a CPU e a GPU e não é necessário transferir dados). Uma das características de uma placa de vídeo é a largura de banda da memória ou largura de banda da memória, que determina a largura de banda teórica da placa. Para o Tesla k80, é de 480 GB / s, para o Tesla v100, já é de 900 GB / s. Além disso, a versão PCI Express e a implementação de como você transferirá dados para o cartão afetarão a taxa de transferência, por exemplo, isso pode ser feito em vários fluxos paralelos.

Vejamos os resultados práticos obtidos para a placa gráfica Tesla k80 na nuvem Amazon:


Hora de transferir dados para a GPU, classificar e transferir dados de volta para a RAM em ms

HtoD - transferir dados para a placa de vídeo da

GPU Execução - classificar na placa de vídeo

DtoH - copiar dados da placa de vídeo para a RAM

A primeira coisa a observar é que a leitura de dados da placa de vídeo é mais rápida do que escreva-as lá em baixo.

A segunda - ao trabalhar com uma placa de vídeo, é possível obter latência de 350 microssegundos, e isso já pode ser suficiente para alguns aplicativos de baixa latência.

O gráfico abaixo mostra uma sobrecarga para mais dados:


Hora de transferir dados para a GPU, classificar e transferir dados de volta para a RAM em ms

Uso do servidor


A pergunta mais comum é como a placa de vídeo do jogo difere da placa de servidor? De acordo com as características, eles são muito semelhantes, mas os preços diferem significativamente.


As principais diferenças entre o servidor (NVIDIA) e a placa de jogo:

  • Garantia do fabricante (o cartão de jogo não foi projetado para uso do servidor)
  • Possíveis problemas de virtualização para uma placa gráfica de consumidor
  • Disponibilidade do mecanismo de correção de erros na placa do servidor
  • O número de threads paralelos (não núcleos CUDA) ou suporte ao Hyper-Q, que permite trabalhar com o cartão de vários threads na CPU, por exemplo, fazer upload de dados para o cartão de um thread e iniciar cálculos com outro

Essas são, talvez, as principais diferenças importantes que encontrei.

Multithreading


Depois que descobrimos como executar o algoritmo mais simples na placa de vídeo e quais resultados podem ser esperados, a próxima pergunta lógica é como a placa de vídeo se comportará ao processar várias solicitações paralelas. Como resposta, tenho dois gráficos de computação na GPU e um processador com 4 e 32 núcleos:


O tempo necessário para executar cálculos matemáticos na GPU e na CPU com matrizes de 1000 x 60 em ms.Este

gráfico executa cálculos com matrizes de 1000 x 60 elementos. Os cálculos são iniciados a partir de vários fluxos de programas; um fluxo separado é criado para a GPU para cada fluxo de CPU (o próprio Hyper-Q é usado). 

Como você pode ver, o processador lida muito bem com essa carga, enquanto a latência de uma solicitação por GPU aumenta significativamente com um aumento no número de solicitações paralelas.


O tempo para realizar cálculos matemáticos na GPU e CPU com matrizes 10.000 x 60 em ms.

No segundo gráfico, os mesmos cálculos, mas com matrizes 10 vezes mais longas, e a GPU se comporta muito melhor sob essa carga. Esses gráficos são muito indicativos e podemos concluir: o comportamento sob carga depende da natureza da própria carga. Um processador também pode lidar com cálculos matriciais de maneira bastante eficiente, mas até certo ponto. Para uma placa de vídeo, é característico que, para uma pequena carga de computação, o desempenho diminua aproximadamente linearmente. Com um aumento na carga e no número de threads paralelos, a placa de vídeo lida melhor. 

É difícil supor como a GPU se comportará em várias situações, mas como você pode ver, sob certas condições, uma placa de servidor pode processar solicitações de vários fluxos paralelos com bastante eficiência.

Discutiremos mais algumas perguntas que você possa ter se ainda decidir usar a GPU em seus projetos.

Limite de recursos


Como já dissemos, os dois principais recursos de uma placa de vídeo são os núcleos e a memória da computação.

Por exemplo, temos vários processos ou contêineres usando uma placa de vídeo e gostaríamos de poder compartilhar a placa de vídeo entre eles. Infelizmente, não existe uma API simples para isso. A NVIDIA oferece tecnologia vGPU , mas não encontrei a placa Tesla k80 na lista de placas suportadas e, até onde posso entender pela descrição, a tecnologia está mais focada em monitores virtuais do que em cálculos. Talvez a AMD ofereça algo mais adequado.

Portanto, se você planeja usar a GPU em seus projetos, deve confiar no fato de que o aplicativo usará a placa de vídeo exclusivamente ou controlará programaticamente a quantidade de memória alocada e o número de núcleos usados ​​nos cálculos.

Contêineres e GPU


Se você descobriu o limite de recursos, a seguinte pergunta lógica: e se houver várias placas de vídeo no servidor?

Novamente, você pode decidir, no nível do aplicativo, qual GPU ele usará.

Outra maneira mais conveniente são os contêineres Docker. Você pode usar contêineres regulares, mas a NVIDIA oferece seus contêineres NGC , com versões otimizadas de vários softwares, bibliotecas e drivers. Para um contêiner, você pode limitar o número de GPUs usadas e sua visibilidade para o contêiner. A sobrecarga no uso de contêiner é de cerca de 3%.

Trabalhar em um cluster


Outra pergunta, o que fazer se você deseja executar uma tarefa em várias GPUs no mesmo servidor ou cluster?

Se você escolheu uma biblioteca semelhante ao empuxo ou a uma solução de nível inferior, a tarefa terá que ser resolvida manualmente. Estruturas de alto nível, por exemplo, para aprendizado de máquina ou redes neurais, geralmente suportam a capacidade de usar várias placas prontas para uso.

Além disso, gostaria de observar que, por exemplo, a NVIDIA oferece uma interface para troca direta de dados entre cartões - NVLINK , que é significativamente mais rápido que o PCI Express. E há tecnologia para acesso direto à memória da placa de outros dispositivos PCI Express - GPUDirect RDMA , incl. e rede .

Recomendações


Se você está pensando em usar a GPU em seus projetos, é provável que a GPU seja adequada para você se:

  • Sua tarefa pode ser reduzida para uma visualização SIMD
  • É possível carregar a maioria dos dados no mapa antes dos cálculos (cache)
  • O desafio envolve computação intensiva

Você também deve fazer perguntas com antecedência:

  • Quantas consultas paralelas serão 
  • Que latência você espera
  • Você precisa de uma placa para sua carga? Você precisa de um servidor com várias placas ou um cluster de servidores GPU 

Só isso, espero que o material seja útil e ajude você a tomar a decisão certa!

Referências


Referência e resultados no github - https://github.com/tishden/gpu_benchmark/tree/master/cuda

Além do tópico, uma gravação do relatório "Bancos de dados GPU - Arquitetura, desempenho e perspectivas de uso"

Webinar NVIDIA NGC Containers - http : //bit.ly/2UmVIVt ou http://bit.ly/2x4vJKF

More articles:


All Articles