Computação paralela com CUDA (Parte 2)
Anteriormente falamos sobre computação paralela e de como ela se diferencia da computação serial. Nesse texto desenvolveremos um código simples para demonstrar como se dá a codificação e execução de um programa paralelo em uma arquitetura heterogênea, empregando a CPU, GPU e o conjunto de interfaces CUDA. Para tanto, espera-se que o leitor disponha de uma placa de vídeo NVIDIA em seu computador, com o kit de desenvolvimento CUDA devidamente instalado e configurado — aqui consta um bom tutorial sobre como instalar e configurar CUDA no Windows. E embora o conhecimento prévio da linguagem C não seja um requisito necessário, os leitores com algum conhecimento de C, sobretudo com o uso de ponteiros e alocação dinâmica de memória, poderão se beneficiar disso.
Todo programa CUDA apresenta uma estrutura básica que é diretamente relacionada ao ambiente heterogêneo de programação. Este ambiente é composto pela CPU e GPU, cada uma com suas respectivas memórias, além de uma interface de comunicação entre ambas. Em decorrência disso, o programa consiste em duas porções de código (1) host code, a ser executado pela CPU fazendo uso de sua memória e (2) device code, a ser executado na GPU utilizando também sua memória dedicada. Observe que essas porções de código não são postas disjuntas, pelo contrário, as linhas são escritas de forma a viabilizar a maximização da eficiência, executando os códigos de forma alternada.
Sendo assim, conforme vimos anteriormente, a depender do problema para o qual estamos programando a solução, o código será escrito de forma serial para a maior parte da solução, transferindo o fluxo de execução para a GPU sempre que o compilador identificar a linha que indica essa ação. Como na figura seguinte.
A boa notícia é que CUDA simplifica bastante o uso da GPU para fazer cálculos, sobretudo quando a natureza do problema permite a paralelização. Para esse texto, faremos a operação mais simples possível, a soma de dois vetores. Embora simples, esse problema nos permite enxergar a dinâmica de programação de um problema utilizando CUDA, aliás, a dinâmica completa, qual seja: (1) carregar os dados na memória do computador, (2) transferir os dados para a memória da GPU, (3) executar os cálculos paralelos na GPU e (4) transferir os resultados dos cálculos de volta para o host, liberando a memória anteriormente alocada na CPU. Caso o leitor não esteja familiarizado com esse processo recomendamos uma visita a parte I dessa publicação.
Nesse primeiro exemplo prático sobre programação paralela com CUDA vamos criar um programa que é o típico Hello World referente a programação CUDA, a soma de dois vetores. A estratégia que seguiremos será, a partir do código C para executar essa soma, adaptá-lo para que essa operação seja feita pela GPU.
Aqui escrevemos a função somaVetor() e definimos como parâmetros os ponteiros A, B e C e uma variável do tipo inteiro para definir o tamanho dos arrays. Na função, muito simples, utilizamos um laço para fazer a soma componente-a-componente.
No corpo do programa fazemos a chamada da função.
Aqui alocamos o espaço necessário, no primeiro laço nós preenchemos os vetores, depois a função é executada e por fim fazemos a impressão do resultado, liberando a memória. Agora vamos adaptar o código para executá-lo na GPU.
Chamamos atenção para o fato de que uma função a ser executada na GPU é denominada kernel e para transformar a função somaVetor em um kernel são necessárias algumas adaptações.
O identificador __global__ indica que essa função pode ser chamada pela CPU, mas que é executada na GPU. Mudamos o nome da função para addKernel apenas para estabelecermos essa diferença. Os parâmetros são os mesmos, com exceção para o fato de retirarmos o parâmetro indicador do tamanho do array.
Cada thread dentro de um bloco possui um índice próprio, um endereço que pode ser acessado utilizando a função threadIdx.x.
Se pensarmos no bloco como um retângulo, cada linha do bloco é formada por threads. Cada thread na primeira linha terá coordenada y (threadIdx.y=0) e coordenada x (threadIdx.x entre 0 e 6). Na segunda linha a coordenada y será igual a 1 e assim sucessivamente. Em nosso exemplo, o bloco possui apenas uma linha, portanto, só precisamos lidar com a coordenada x dos threads.
O que a segunda linha da função faz é justamente isso, cada thread então fará a soma (em paralelo) de uma linha dos vetores, com o resultado sendo alocado em c.
Agora nós alocamos a memória necessária no host e preenchemos os vetores com valores hipotéticos. Por fim, declaramos três ponteiros os quais serão utilizados para armazenar o endereço da memória a ser alocada no dispositivo, ou seja, na GPU. Em geral, os ponteiros que apontam para posições na memória da GPU são nomeados utilizando esse padrão: dev_nomedavariávelnohost. Isso deixa claro para quem lê o código a natureza dessas variáveis.
Nessa parte do código é executada a alocação de memória na GPU e os respectivos endereços são armazenados nos ponteiros anteriormente declarados (dev_a, dev_b, dev_c). A função cudaMalloc() é responsável pela alocação dinâmica de memória do processador gráfico. Na chamada da função, nós passamos os ponteiros por referência (por isso os ponteiros duplos em void**) e definimos o tamanho da memória requerida.
Agora nós copiamos os vetores a e b para a memória da GPU (lembre-se que a GPU só acessa sua própria memória) usando a função cudaMemcpy(destino, origem, tamanho, Origem-Destino). Após isso, é executada a chamada da função kernel.
Sempre que ocorre a chamada do kernel, são utilizados os símbolos <<<númeroDeBlocos, threadsPorBlocos>>>. Aqui nós definimos o número de threads de modo que cada um deles somará uma linha dos vetores (pense no vetor coluna), para tanto, definimos apenas 1 bloco (com uma dimensão) com tamanho (size=6) threads. Note também que os parâmetros da função são os vetores apontando para a memória da GPU, não do host.
Depois que o kernel é executado, nós copiamos o resultado de volta para a memória do host utilizando a função cudaMemcpy(destino, origem, tamanho, origem-destino). Em seguida, liberamos a memória da GPU e imprimimos o resultado em c, liberando a memória do heap.
Esse é o resultado final do programa. A operação com mais alta latência em todo esse programa é a transferência de dados entre dispositivos e quanto maior for o volume de dados, mais tempo o programa leva para ser executado, sobretudo no caso em que haja muitas transferências de dados. Felizmente, em muitos casos, pelo menos no que tange às simulações, existe uma biblioteca para geração de número aleatórios (cuRAND) e também para operações de álgebra linear (cuBLAS) com funções executadas diretamente na GPU. No caso de uma simulação gerada a partir de números aleatórios, os dados podem ser gerados diretamente na memória da GPU e apenas os resultados finais são transferidos pra o host, minimizando a latência das operações.
Esse foi um exemplo básico de como escrever um programa em CUDA C, nele executamos os processos básicas de transferência de dados da memória do host para a memória do processador gráfico, execução do programa em paralelo pela GPU e posterior transferência ao host dos resultados obtidos na GPU.
Em um próximo texto, os threads serão analisados com maior nível de detalhes. Veremos que threads são estruturados em blocos e que blocos são estruturados em grades. Analisando esses pormenores, veremos como essa estrutura faz sentido e facilita o trabalho de manipulação de grandes volumes de dados em programação paralela. Ao trabalharmos com CUDA, compreender bem essa estrutura é essencial para implementação de soluções envolvendo problemas complexos.
Abaixo disponibilizamos o código completo.
