2 dias brincando com Programação na GPU

Hey meus amigos, neste final de semana comprei meu novo computador pessoal, pela primeira vez adquiri um computador que possuia uma placa de vídeo NVidia, nunca fui tão apaixonado por jogos, porém sempre gostei muito de programação concorrente e distribuida, e queria experimentar o potencial do novo brinquedo.

Conhecia o projeto CUDA da NVidia, que permite que desenvolvedores possam escrever programas e serem executados pela própria GPU, como pythonista também encontrei o projeto Numba que permite com que código python seja convertido em código CUDA via LLVM.

O Conceito

Por que utilizar a GPU ?

A GPU é potencialmente desenhada para realizar grandes blocos de processamento em paralello, como jogos, vídeos e imagens, que exigem uma performance crucial para a experiência.

Devido a essa performance e paralelismo, ela é utilizada também para processar grandes massas de dados necessárias para machine learn, big data, mineração de bitcoins e outras possibilidades.

O que muda na cabeça do desenvolvedor para programar utilizando a GPU ?

Primeiro, o desafio é pensar na possibilidade de quebrar o algoritmo em vários steps que possa ser executado em paralelo, imagine como poderiarmos separar centenas ou milhares de mentes para resolver um problema o mais eficiente possível ?, Imagine orquestrar 5760 cores de processamento utilizando uma GeForce Gtx Titan Z por exemplo, claro que não comprei uma placa tão cara como esta, peguei um modelo que possui 384 cores de processamento.

Começando a brincar

Instalação dos brinquedos

• Python 3.7
• Cuda Toolkit: https://developer.nvidia.com/cuda-downloads, como utilizo archlinux, instalei via: pacman -S cuda python-pycuda
• Numba: pip install numba

Entender como funciona o paralelismo do CUDA

1. Kernels
   São funções que são levadas para o dispositivo e podem executados N vezes por diferentes cores da GPU, no exemplo abaixo o numba compila a função utilizando LLVM e leva para GPU, bastando decorar a função com @cuda.jit .

neste exemplo declaramos uma função que multiplicará uma posição dentro da matrix por 2.

2. Disparar blocos e threads para executar esta função
Beleza criamos nossa primeira função que faz uma operação bem simples, agora precisamos dizer para a GPU rodar essa mesma função para uma massa de dados.

O resultado esperado seria [0, 2, 4, … 1998]

Temos um conceito bem interresante que o CUDA organiza as execuções utilizando blockspergrid e threadsperblock , assim cria formas de executar os kernels por blocos e threads, o conceito pode ser explicado amplamente aqui: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#introduction.

3. Executando
No archlinux precisei dizer para a aplicação aonde estava localizado as biblitecas do CUDA.
Executei da seguinte forma.

env NUMBAPRO_NVVM=/opt/cuda/nvvm/lib64/libnvvm.so NUMBAPRO_LIBDEVICE=/opt/cuda/nvvm/libdevice python arquivo.py

Memória da máquina x memória da GPU

Imagine que precisaremos rodar a função add_by_2 e o resultado dela séra input de outra função, como podemos otimizar o repasse de variáveis ? vamos lá demonstrar.

Cada vez que disparamos uma execução, o numba copia os parametros que estão na memória principal para a memória da GPU, e acabando a execução copia da GPU para a memória da máquina, podemos explicitamente levar o array para a GPU e recuperar quando acreditarmos que é necessário, olhe o exemplo abaixo.

Esta função gastou 2 segundos para ser executada

Podemos levar o conteúdo de an_array para a GPU utilizando cuda_array = cuda.to_device(array) e recuperar utilizando cuda_array.copy_to_host()

Gastamos 500ms na execução acima, só recuperando da memória no momento certo.

Tentativa de implementar um algoritmo de sort totalmente paralello

Depois de facinar pelo brinquedo, queria um desafio para aprender mais das entranhas do CUDA e GPU Programming.

Exemplo de sort acontecendo na GPU

Li o artigo sobre como executar um algoritmo de sort dentro da GPU: https://www.alanzucconi.com/2017/12/13/gpu-sorting-1/, e comecei a projetar um rascunho.

Exemplo de algoritmo utilizando Odd-Even Sort

Minha primeira implementação de um algoritmo de sort utilizando GPU

Observações

• Para cada par dentro do vetor eu disparo uma thread de ordenação, que ela realiza swaps tanto a esquerda quanto a direita para levar os números para as posições corretas.
• O algortimo ainda possui um race-condition que ainda preciso estudar, alguns números ficam duplicados e sobrescrevem algumas posições.
• A performance cai drasticamente quando vamos para um milhão de números, na linha quatorze fazemos um loop de N/2, e são N/2 execuções paralelas, talvez eu precise de um algoritmo mais eficiente.

Possibilidades

• A GPU possui uma variedades de instruções matemáticas fantásticas: tudo pronto e só utilizar: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html

Conclusão

Nestes dois dias brinquei e estudei um horizonte totalmente novo para mim, acredito que se eu for me aprofundar no assunto necessitarei estudar e praticar bastante, como relembrar algebra linear, funções matemáticas, estudar mais algoritmos, computação sempre nos desafia a aprender mais e realizar tarefas com maior eficiencia.

Referências

Numba for CUDA: http://numba.pydata.org/numba-doc/dev/cuda/
Cuda programming guide: https://docs.nvidia.com/cuda/cuda-c-programming-guide/
GPU Sorting: https://www.alanzucconi.com/2017/12/13/gpu-sorting-1/
Meu repositório de playground com CUDA: https://github.com/wpjunior/cuda-numba-playground