Que tipos de problemas se prestam bem à computação em GPU?

84

Portanto, tenho uma cabeça decente para quais problemas com os quais trabalho são os melhores em série e que podem ser gerenciados em paralelo. Mas, no momento, não tenho muita idéia do que é melhor manipulado pela computação baseada em CPU e do que deve ser transferido para uma GPU.

Eu sei que é uma pergunta básica, mas grande parte da minha pesquisa é capturada por pessoas que claramente defendem uma ou outra sem realmente justificar o porquê , ou regras práticas um tanto vagas. Procurando uma resposta mais útil aqui.

Fomite
fonte

Respostas:

63

O hardware da GPU possui dois pontos fortes: FLOPs (computação bruta) e largura de banda da memória. Os problemas computacionais mais difíceis se enquadram em uma dessas duas categorias. Por exemplo, álgebra linear densa (A * B = C ou Resolver [Ax = y] ou Diagonalizar [A], etc) cai em algum lugar no espectro da largura de banda de computação / memória, dependendo do tamanho do sistema. As transformadas rápidas de Fourier (FFT) também se encaixam nesse molde com altas necessidades de largura de banda agregada. Assim como outras transformações, algoritmos baseados em grade / malha, Monte Carlo, etc. Se você observar os exemplos de código do NVIDIA SDK , poderá ter uma idéia dos tipos de problemas mais comuns.

Acho que a resposta mais instrutiva é a pergunta "Em que tipos de problemas as GPUs são realmente ruins?" A maioria dos problemas que não se enquadram nessa categoria pode ser executada na GPU, embora alguns exijam mais esforço do que outros.

Problemas que não são bem mapeados geralmente são muito pequenos ou imprevisíveis. Problemas muito pequenos carecem do paralelismo necessário para usar todos os threads na GPU e / ou podem caber em um cache de baixo nível na CPU, aumentando substancialmente o desempenho da CPU. Problemas imprevisíveis têm muitas ramificações significativas, o que pode impedir que os dados fluam eficientemente da memória da GPU para os núcleos ou reduza o paralelismo quebrando o paradigma SIMD (consulte ' warps divergentes '). Exemplos desses tipos de problemas incluem:

  • A maioria dos algoritmos de gráficos (imprevisível demais, especialmente no espaço da memória)
  • Álgebra linear esparsa (mas isso também é ruim na CPU)
  • Pequenos problemas de processamento de sinal (FFTs menores que 1000 pontos, por exemplo)
  • Procurar
  • Ordenar
Max Hutchinson
fonte
3
Ainda assim, as soluções de GPU para esses problemas "imprevisíveis" são possíveis e, embora hoje em dia não sejam normalmente viáveis, podem ganhar significado no futuro.
precisa saber é o seguinte
6
Eu gostaria de adicionar especificamente ramificações à lista de disjuntores de desempenho da GPU. Você deseja que todas as suas (centenas) executem a mesma instrução (como no SIMD) para executar uma computação verdadeiramente paralela. Por exemplo, nas placas AMD, se algum dos fluxos de instruções encontrar uma ramificação e precisar divergir - toda a frente de onda (grupo paralelo) diverge. Se outras unidades da frente da onda não divergirem - elas devem executar uma segunda passagem. Isso é o que maxhutch quer dizer com previsibilidade, eu acho.
Girafa violeta
2
@VioletGiraffe, isso não é necessariamente verdade. No CUDA (ou seja, nas GPUs da Nvidia), a divergência de ramificação afeta apenas o warp atual, que possui no máximo 32 threads. Warps diferentes, embora executando o mesmo código, não são síncronos, a menos que explicitamente sincronizados (por exemplo, com __synchtreads()).
Pedro
1
@ Pedro: É verdade, mas a ramificação em geral prejudica o desempenho. Para códigos de alto desempenho (que código de GPU não é?), É quase essencial levar isso em conta.
jvriesem
21

Os problemas com alta intensidade aritmética e padrões regulares de acesso à memória são geralmente fáceis de implementar nas GPUs e têm um bom desempenho.

A dificuldade básica de se ter um código de GPU de alto desempenho é que você tem uma tonelada de núcleos e deseja que todos sejam utilizados com a máxima potência possível. Problemas que possuem padrões irregulares de acesso à memória ou que não possuem alta intensidade aritmética dificultam as coisas: você passa muito tempo comunicando resultados ou muito tempo buscando coisas da memória (que é lenta!), E não tempo suficiente analisando números. É claro que o potencial de simultaneidade no seu código é fundamental para a capacidade de ser implementado também na GPU.

Reid.Atcheson
fonte
Você pode especificar o que quer dizer com padrões regulares de acesso à memória?
fomite
1
A resposta de maxhutch é melhor que a minha. O que quero dizer com padrão de acesso regular é que a memória é acessada de maneira temporal e espacialmente localmente. Ou seja: você não faz saltos enormes em torno da memória repetidamente. Também é uma espécie de pacote que eu notei. Também se entende que seus padrões de acesso a dados podem ser predeterminados de alguma forma pelo compilador ou pelo programador, para que a ramificação (instruções condicionais no código) seja minimizada.
usar o seguinte código
15

Isso não pretende ser uma resposta por si só, mas sim uma adição às outras respostas de maxhutch e Reid.Atcheson .

Para tirar o melhor proveito das GPUs, seu problema não precisa apenas ser altamente (ou maciçamente) paralelo, mas também o algoritmo principal que será executado na GPU, deve ser o menor possível. Nos termos do OpenCL , isso é chamado principalmente de kernel .

Para ser mais preciso, o kernel deve caber no registro de cada unidade de multiprocessamento (ou unidade de computação ) da GPU. O tamanho exato do registro depende da GPU.

Dado que o kernel é pequeno o suficiente, os dados brutos do problema precisam se encaixar na memória local da GPU (leia-se: memória local (OpenCL) ou memória compartilhada (CUDA) de uma unidade de computação). Caso contrário, mesmo a grande largura de banda de memória da GPU não é rápida o suficiente para manter os elementos de processamento ocupados o tempo todo.
Geralmente, essa memória tem cerca de 16 a 32 KiByte de tamanho .

Torbjörn
fonte
A memória local / compartilhada de cada unidade de processamento não é compartilhada entre todas as dezenas (?) De threads em execução em um único cluster de núcleos? Nesse caso, você não precisa realmente manter seu conjunto de dados significativamente menor para obter o desempenho total da GPU?
Dan Neely
A memória local / compartilhada de uma unidade de processamento é acessível apenas pela própria unidade de computação e, portanto, compartilhada apenas pelos elementos de processamento dessa unidade de computação. A memória global da placa gráfica (normalmente 1 GB) é acessível por todas as unidades de processamento. A largura de banda entre os elementos de processamento e a memória local / compartilhada é muito rápida (> 1 TB / s), mas a largura de banda da memória global é muito mais lenta (~ 100 GB / s) e precisa ser compartilhada entre todas as unidades de computação.
Torbjörn 23/01
Eu não estava perguntando sobre a principal memória da GPU. Eu pensei que a memória on die fosse alocada apenas no cluster do nível do núcleo, não por núcleo individual. ex para um nVidia GF100 / 110 gpu; para cada um dos 16 clusters SM, não os 512 núcleos cuda. Com cada SM projetado para executar até 32 threads em paralelo, maximizar o desempenho da GPU exigiria manter o conjunto de trabalho na faixa de 1kb / thread.
Dan Neely
@Torbjoern O que você deseja é manter todos os pipelines de execução da GPU ocupados, as GPUs conseguem isso de duas maneiras: (1) a maneira mais comum é aumentar a ocupação, ou dito de forma diferente, aumentando o número de threads simultâneos (kernels pequenos usam menos de os recursos compartilhados para que você possa ter threads mais ativos); talvez o melhor seja (2) aumentar o paralelismo do nível de instrução dentro do seu kernel, para que você possa ter um kernel maior com ocupação relativamente baixa (pequeno número de threads ativos). Veja bit.ly/Q3KdI0
fcruz
11

Provavelmente, uma adição mais técnica às respostas anteriores: GPUs CUDA (Nvidia) podem ser descritas como um conjunto de processadores que funcionam autonomamente em 32 threads cada. Os threads em cada processador funcionam na etapa de bloqueio (pense no SIMD com vetores de comprimento 32).

Embora a maneira mais tentadora de trabalhar com GPUs seja fingir que absolutamente tudo funciona em um travamento, essa nem sempre é a maneira mais eficiente de fazer as coisas.

Se o seu código não não paralelizar bem / automaticamente para centenas / milhares de tópicos, você pode ser capaz de dividi-la em tarefas assíncronas individuais que não paralelizar bem, e executar aqueles com apenas 32 threads em execução em lock-passo. O CUDA fornece um conjunto de instruções atômicas que possibilitam implementar mutexes que, por sua vez, permitem que os processadores sincronizem entre si e processem uma lista de tarefas em um paradigma de conjunto de encadeamentos . Seu código funcionaria da mesma maneira que em um sistema com vários núcleos, mas lembre-se de que cada núcleo possui 32 threads.

Aqui está um pequeno exemplo, usando CUDA, de como isso funciona

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Você precisa chamar o kernel main<<<N,32>>>(tasks,nr_tasks)para garantir que cada bloco contenha apenas 32 threads e, portanto, se encaixe em um único warp. Neste exemplo, também assumi, por simplicidade, que as tarefas não têm nenhuma dependência (por exemplo, uma tarefa depende dos resultados de outra) ou conflitos (por exemplo, trabalham na mesma memória global). Se for esse o caso, a seleção de tarefas se tornará um pouco mais complicada, mas a estrutura é essencialmente a mesma.

Obviamente, isso é mais complicado do que apenas fazer tudo em um grande lote de células, mas amplia significativamente o tipo de problemas para os quais as GPUs podem ser usadas.

Pedro
fonte
2
Isso é tecnicamente verdadeiro, mas é necessário alto paralelismo para obter alta largura de banda de memória e há um limite para o número de chamadas assíncronas do kernel (atualmente 16). Há também toneladas de comportamento não documentado relacionado ao agendamento na versão atual. Eu aconselho contando com kernels assíncronos para imporove desempenho para o momento ...
Max Hutchinson
2
O que estou descrevendo pode ser feito em uma única chamada do kernel. Você pode fazer N blocos de 32 threads cada, de modo que cada bloco se encaixe em um único warp. Cada bloco adquire uma tarefa de uma lista global de tarefas (acesso controlado usando átomos / mutexes) e calcula-a usando 32 threads com etapas de bloqueio. Tudo isso acontece em uma única chamada do kernel. Se você quiser um exemplo de código, entre em contato e postarei um.
Pedro
4

Um ponto não mencionado até agora é que a geração atual de GPUs não se sai tão bem em cálculos de ponto flutuante de precisão dupla quanto em cálculos de precisão única. Se seus cálculos precisarem ser feitos com precisão dupla, você poderá esperar que o tempo de execução aumente em um fator de 10 ou mais sobre a precisão única.

Brian Borchers
fonte
Eu quero discordar. A maioria (ou todas) as GPUs mais recentes têm suporte a precisão dupla nativa. Quase todas essas GPU relatam cálculos de precisão dupla, executando aproximadamente a metade da velocidade da precisão única, provavelmente devido à duplicação simples dos acessos à memória / largura de banda necessários.
Godric Seer
1
Embora seja verdade que as melhores e mais recentes placas Nvidia Tesla oferecem desempenho de pico de precisão dupla, que é metade do desempenho de pico de precisão única, a proporção é de 8 para 1 para os cartões de consumo da arquitetura Fermi mais comuns.
precisa
@GodricSeer A proporção 2: 1 de ponto flutuante SP e DP tem muito pouco a ver com largura de banda e quase tudo a ver com quantas unidades de hardware existem para executar essas operações. É comum reutilizar o arquivo de registro para SP e DP, portanto, a unidade de ponto flutuante pode executar 2x as operações de SP como operações de DP. Existem inúmeras exceções nesse design, por exemplo, IBM Blue Gene / Q (não possui lógica SP e, portanto, o SP é executado em ~ 1,05x DP). Alguns GPUs têm outras de 2 proporções, por exemplo, 3 e 5.
Jeff
Faz quatro anos que escrevi essa resposta, e a situação atual das GPUs da NVIDIA é que, para as linhas GeForce e Quadro, a proporção DP / SP agora é de 1/32. As GPUs Tesla da NVIDIA têm desempenho de precisão dupla muito mais forte, mas também custam muito mais. Por outro lado, a AMD não prejudicou o desempenho de precisão dupla em suas GPUs Radeon da mesma maneira.
Brian Borchers
4

Do ponto de vista metafórico, a gpu pode ser vista como uma pessoa deitada em uma cama de unhas. A pessoa que está no topo são os dados e, na base de cada unha, existe um processador; portanto, a unha é na verdade uma seta apontando do processador para a memória. Todas as unhas estão em um padrão regular, como uma grade. Se o corpo está bem espalhado, é bom (o desempenho é bom); se o corpo toca apenas alguns pontos do leito ungueal, então a dor é ruim (mau desempenho).

Isso pode ser tomado como uma resposta complementar às excelentes respostas acima.

labotsirc
fonte
4

Pergunta antiga, mas acho que essa resposta de 2014 - relacionada a métodos estatísticos, mas generalizável para quem sabe o que é um loop - é particularmente ilustrativa e informativa.

GT.
fonte
2

As GPUs têm E / S de longa latência, portanto, muitos threads precisam ser usados ​​para saturar a memória. Manter um warp ocupado requer muitos threads. Se o caminho do código for 10 e a latência de E / S 320, 32 threads deverão aproximar-se da saturação do warp. Se o caminho do código for 5 relógios, dobre os threads.

Com mil núcleos, procure milhares de threads para utilizar totalmente a GPU.

O acesso à memória é por linha de cache, geralmente 32 bytes. Carregar um byte tem um custo comparável a 32 bytes. Portanto, combine o armazenamento para aumentar a localidade de uso.

Existem muitos registros e RAM local para cada warp, permitindo o compartilhamento de vizinhos.

Simulações de proximidade de grandes conjuntos devem otimizar bem.

E / S aleatória e segmentação única são uma alegria para matar ...

user14381
fonte
Esta é uma pergunta genuinamente fascinante; Estou discutindo comigo mesmo se é possível (ou vale o esforço) 'paralelizar' uma tarefa razoavelmente direta (detecção de borda em imagens aéreas) quando cada tarefa leva ~ 0,06s, mas existem ~ 1,8 milhões de tarefas a serem executadas ( por ano, durante 6 anos de dados: as tarefas são definitivamente separáveis) ... portanto, aproximadamente 7,5 dias de tempo de computação em um único núcleo. Se cada cálculo fosse mais rápido em uma GPU e o trabalho pudesse ser paralelo 1 por nGPUcores [n pequeno], é realmente provável que o tempo de trabalho caia para ~ 1 hora? Parece improvável.
GT.
0

Imagine um problema que possa ser resolvido com muita força bruta, como o Travelling Salesman. Imagine que você tenha racks de servidores com 8 placas de vídeo espancadas cada uma e cada placa tenha 3000 núcleos CUDA.

Simplesmente resolva TODAS as rotas possíveis do vendedor e depois classifique por tempo / distância / alguma métrica. Claro que você está jogando fora quase 100% do seu trabalho, mas a força bruta às vezes é uma solução viável.

Criggie
fonte
Eu tive acesso a um pequeno farm de quatro desses servidores por uma semana e, em cinco dias, fiz mais blocos de distribut.net do que nos 10 anos anteriores.
Criggie
-1

Ao estudar muitas idéias de engenharia, eu diria que uma gpu é uma forma de focar tarefas, de gerenciamento de memória e de cálculo repetitivo.

Muitas fórmulas podem ser simples de escrever, mas difíceis de calcular, como na matemática matricial, você não recebe uma única resposta, mas muitos valores.

Isso é importante na computação, com a rapidez com que um computador calcula valores e executa fórmulas, pois algumas fórmulas não podem ser executadas sem todos os valores calculados (por isso, diminuem a velocidade). Um computador não sabe muito bem qual ordem executar fórmulas ou calcular valores a serem usados ​​nesses programas. Principalmente, as forças brutas passam a altas velocidades e dividem as fórmulas em mandris para calcular, mas muitos programas hoje em dia exigem esses mandris calculados agora e aguardam perguntas (e questões e mais e mais).

Por exemplo, em um jogo de simulação que deve ser calculado primeiro em colisões, o dano da colisão, a posição dos objetos, a nova velocidade? Quanto tempo isso deve levar? Como qualquer CPU pode lidar com essa carga? Além disso, a maioria dos programas é muito abstrata, exigindo mais tempo para manipular dados e nem sempre é projetada para multiencadeamento ou não há boas maneiras em programas abstratos para fazer isso de maneira eficaz.

À medida que a cpu se tornou cada vez melhor, as pessoas ficaram desleixadas na programação e também precisamos programar para muitos tipos diferentes de computadores. Uma gpu é projetada para fornecer força bruta através de muitos cálculos simples ao mesmo tempo (sem mencionar a memória (secundária / ram) e o resfriamento por aquecimento são os principais gargalos da computação). Uma cpu está gerenciando muitas questões ao mesmo tempo ou sendo puxada para várias direções e está tentando descobrir o que fazer para não conseguir. (ei, é quase humano)

Um gpu é um trabalhador pesado, o trabalho tedioso. Uma CPU está gerenciando o caos completo e não pode lidar com todos os detalhes.

Então o que aprendemos? Uma gpu detalha o trabalho tedioso de uma só vez e uma cpu é uma máquina de múltiplas tarefas que não consegue se concentrar muito bem com muitas tarefas a serem executadas. (É como se tivesse desordem de atenção e autismo ao mesmo tempo).

Engenharia existem as idéias, design, realidade e muito trabalho pesado.

Ao sair, lembre-se de começar simples, comece rapidamente, falhe rapidamente, falhe rapidamente e nunca pare de tentar.

Andrew G. Corbi
fonte