Vitor Pamplona

Innovation on Vision: Imaging , Enhancement and Simulation

Cuda

O hardware gráfico é famoso pelo alto poder computacional ao processar streams de dados. Até meados de 2004, estes streams eram apenas conjuntos de vértices, pontos ou pixels de jogos e ferramentas de visualização. Com a linguagem Cuda a nVidia tenta expandir seu público alvo, possibilitando o uso do hardware gráfico por outros desenvolvedores.

 Cuda (Compute Unified Device Architecture) é uma linguagem para programação paralela de propósito geral desenvolvido pela nVidia para placas gráficas da nVidia (Séries + 8xxx, Tesla, Quadro). Fortemente baseada em C, a linguagem não requer conhecimento do hardware gráfico (GPU), do pipeline da OpenGL ou sequer do uso de Texturas, uma das formas mais utilizadas para trafegar informação entre CPU e GPU. Está disponível gratuitamente, é multi-plataforma (Windows / Linux / Mac), já possui versão estável, é bem documentada e suportada pela nVidia.  

Cuda pode ser considerada uma API de baixo nível, pois expõe os tipos diferentes de memória da placa gráfica e exige que o desenvolvedor configure seus acessos a memória global, a cache e a quantidade e disposição das threads para atingir a melhor performance. Nada é automático. A placa gráfica é vista como um simples dispositivo operando como um co-processador, e aguardando instruções da aplicação cliente. O desenvolvedor é o responsável pelo escalonamento de atividades entre a CPU (Host) e a GPU.

O código que rodará na placa de vídeo é chamado de kernel , um único código que será executado sobre um grande conjunto de informações. As funções de um programa convencional, que são chamadas muitas vezes com entradas e saídas independentes, são as primeiras a se transformarem em kernels.

Por exemplo, uma operação que soma 1 em cada valor de um vetor com 1 GB de dados, pode facilmente ser transformada em um kernel, onde várias threads irão somar 1 e armazenar o valor na mesma posição. Nesse caso, não é necessário sincronia, nem comunicação entre as threads. Perfeito para a GPU.  

A distribuição das threads em Cuda é feita através da definição de blocos no grid da placa gráfica. Numa chamada de método, cada bloco possui o mesmo número de threads. Threads e blocos são referenciados por identificadores 3D. Na primeira imagem, o grid computacional da placa foi criado com 2 linhas e 3 colunas (2,3,1) e cada bloco contém 16 threads (4,4,1). Em tempo de execução, os identificadores - mostrados entre os parênteses - estarão disponíveis, e são eles que definirão a posição na memória global para a busca das informações do processamento. Por exemplo: O nosso vetor de 1 GB de dados é dividido em 6 partes, uma para cada bloco. Por um cálculo simples (blockId.x * 2 + blockId.y), os blocos já sabem qual parte do vetor lhes pertence. Em seguida, divide-se esta parte da memória entre todas as threads do mesmo bloco. Da mesma forma que foi com o grid, as threads também utilizarão um cálculo sobre seus índices para identificarem qual parte da memória do bloco devem processar. Como os endereços no vetor de entrada são os mesmos dos de saída, a paralelização está completa. Claro, o desenvolvedor deve programar tudo isso. Já veremos um exemplo.  

 Cuda possui 6 tipos diferentes de memória como mostrado na figura ao lado. Cada thread possui sua memória local, que são as variáveis declaradas dentro do kernel, e os registradores, que tem a mesma função dos registradores de CPU e não são definidos pelo programador. Cada bloco possui uma memória chamada de memória compartilhada, pois a mesma informação está disponível para todas as threads daquele bloco. Há ainda a memória global, a memória de constantes, e a memória de texturas. As duas últimas são somente leitura e a memória de texturas possui um cache automático.  

Enquanto a memória local, a compartilhada e a memória de constantes são extremamente rápidas de acessar (cerca de 4 flops), a memória global e a memória de textura são muito lentas (cerca de 700 flops). Por esse motivo, é comum no início do processamento, levar trechos mais acessados da memória global para a memória compartilhada e só então executar o algoritmo. O desenvolvedor deve planejar minuciosamente a sua solução, agrupando em blocos as threads que atuam em posições de memória equivalentes ou próximas. Como podem notar, a performance pode variar bastante.  

Cuda não possui stack e a recursão é limitada.  

Vejamos um exemplo de código que copia valores de uma imagem para um outro array de dados. Em CPU esta cópia seria feita sequencialmente, varrendo todos os pixels da imagem pelo mesmo processador. Em GPU, cada processador pode trabalhar em uma parte da imagem, como segue:

// Declara um array bidimensional na memória da placa gráfica. 
cudaArray* cu_array;
texture<float, 2> tex;
 
// Aloca o array na placa gráfica com tamanho widthXheight.
cudaMalloc( &cu_array, cudaCreateChannelDesc<float>(), width, height );
 
// Copia o conteúdo de image que está na CPU para o cu_array que está na GPU.
cudaMemcpy( cu_array, image, width*height, cudaMemcpyHostToDevice);
 
// Bind the array to the texture
cudaBindTexture( tex, cu_array);
 
// Define o número de threads em cada bloco (16x16).
dim3 blockDim(16, 16, 1);

// Determina a quantidade de blocos.
dim3 gridDim(width / blockDim.x, height / blockDim.y, 1);

// Chama a função kernel que já está dentro da GPU
kernel<<< gridDim, blockDim, 0 >>>(d_odata, width, height);

cudaUnbindTexture(tex);

// Esta função será executada na GPU.
// Reparem a forma com que o programador busca o X e o Y para acessar a imagem. 
// A anotação __global__ define que está função pode ser executada na GPU.
__global__ void kernel(float* odata, int height, int width) {
   unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
   unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
   // Busca o valor da posição x, y na textura.
   float c = texfetch(tex, x, y);
   // salva no vetor odata
   odata[y*width+x] = c;
}

Cada thread terá um ID diferente, que se transformará nas variáveis X e Y. Reparem que todos os possíveis valores para X, Y (são valores inteiros) são processados, e são processados uma única vez.

Vale lembrar que a GeForce 9800 GX2 tem 256 processadores de 1.4 Ghz e 1Gb de RAM que ficam ociosos a maioria do tempo em que uma aplicação OpenGL ou DirectX não está rodando. Novas placas tendem a trazer mais processamento. Cabe a nós utilizarmos.

Para quem não gosta de C, já existe o pyCuda (nome legal paiCuda) que simplifica o desenvolvimento. Divirtam-se.

The following pages are talking about "Cuda":

Posted in Dec 14, 2008 by Vitor Pamplona - Edit - History

Showing Comments



- - patrick

- - Posted in May 9, 2009 by 200.203.35.164

Olá Vitor!

Tava pesquisando sobre CUDA, entrei no site da nVidia e vi um monte de tutorial / documentação.

Sabe informar qual daqueles seria o melhor para iniciar os estudos, considerando que não entendo nada do assunto?

Valeu.

- - Max

- - Posted in May 15, 2009 by 189.75.117.66

Oi Max,

Eu usei os tutoriais da Nvidia mesmo. Os exemplos do Toolkit também ajudam a entender.

[] s

- - Posted in May 15, 2009 by Vitor Pamplona

Ótimo Artigo

- - Roussian

- - Posted in Jul 21, 2009 by 189.114.36.215



- - joao junio

- - Posted in Nov 12, 2009 by 201.39.137.216

Olá Vitor,

Gostaria de saber como faço para implementar em CUDA, usando o visual C + +, o seguinte código, feito em C + +:

...
for (k = 0; k < = 50; k + +)
for (j = 0; j < = 50; j + +)
for (i = 0; i < = 50; i + +)
u [i, j, k] = i * j * k;
...

Grato,

Renato

- - Renato

- - Posted in Nov 23, 2009 by 189.114.213.251

Olá, existe um bom site sobre CUDA e GPGPU em português, o endereço é: www.gpubrasil.com. Nele existem vários tutoriais de programação, além de novidades sobre CUDA.

- - Hugo

- - Posted in Feb 7, 2010 by 201.82.131.202

Bacana

- - Sam

- - Posted in Apr 26, 2010 by 189.80.187.75

Olá Vitor. No código que copia valores de uma imagem para um outro array de dados, o seguinte parâmetro image, na chamada:

cudaMemcpy (cu_array, image, width * height, cudaMemcpyHostToDevice),

é realizado usando qual função? opengl?



- - Wallace

- - Posted in Apr 30, 2010 by 189.5.112.247

Add New Comment

Your Name:


Write the code showed above on the text below.