Adição de vetores com CUDA

Vou fazer minha primeira postagem sobre meu trabalho. Isto me ajudou a entender um pouco melhor sobre a tecnologia que estou trabalhando agora (CUDA). Será feita uma adição de 2 vetores. Este pode ser um exemplo "Hello World" para esta tecnologia, pois é um exemplo muito simples e que utiliza da característica paralela do processador.

OBS.: Não é minha intenção explicar os detalhes pequenos do CUDA nem algumas noções básicas, isso pode ser encontrado com muitos detalhes e bem explicado no CUDA programming guide
O CUDA, por ser uma tecnologia de programação de multiprocessadores paralelos em placas de vídeo da NVIDIA tem como característica utilizar muitas threads para poder realizar a computação desejada. Quando fala-se de muitas threads, estamos falando na casa dos milhares. Isto para obter uma eficiência boa com a placa de vídeo. Há outros problemas, como maximizar a quantidade de cálculos efetuados dentro da GPU para que "valha a pena" efetuar esta computação na GPU, no entanto, isto está além do escopo deste artigo.

Bem, mãos à obra.

Inicialmente a primeira coisa a se fazer é criar um arquivo chamado 'adicao_vetores.cu'.
Abre-se então um editor de texto de preferência e começamos a editar o arquivo:
Iniciando como um programa C comum (com algumas coisas básicas do CUDA já incluidas)

#include
#include
int main(int argc, char ** argv) {
CUT_DEVICE_INIT();

return 0;
}

Precisamos de 4 vetores: a, b, resultados_cpu, resultados_gpu (já com valores dentro destes vetores):

int main(int argc, char ** argv) {
CUT_DEVICE_INIT();

int n = 10; //numero de elementos do vetor
float * a = new float[n];
float * b = new float[n];
float * resultados_cpu = new float[n];
float * resultados_gpu = new float[n];

//Dados de entrada
for (int i = 0; i <>void adicionar_vetor_cpu(float * a, float * b, float * resultado, int n) {
for (int i = 0; i <>

Ótimo, agora temos uma soma de vetores implementada em CPU. No entanto, agora temos que adicionar um método para adicionar estes vetores utilizando a GPU:

__global__ void adicionar_vetor(float * a, float * b, float * resultado, int n) {
unsigned int id = threadIdx.x + blockDim.x * blockIdx.x;
if (id <>

• __global__: Esta variável especifica que este é um método de entrada de invocação de computação na GPU. Isto significa que esta função será executada na GPU.
• threadIdx, blockIdx, blockDimx: estas são as variáveis que nos dizem como identificar o ID de uma thread, tornando possível saber onde escrever o resultado da computação

Para executar este método, é preciso alocar memória na GPU para tal e transferir os dados dos vetores de entrada 'a' e 'b' para a GPU. Isto é feito com 'cudaMalloc()' e 'cudaMemcpy':

float * d_a;
float * d_b;
float * d_c;
cudaMalloc((void**)&d_a, sizeof(float) * n);
cudaMalloc((void**)&d_b, sizeof(float) * n);
cudaMalloc((void**)&d_c, sizeof(float) * n);
cudaMemcpy(d_a, a, sizeof(float) * n, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, sizeof(float) * n, cudaMemcpyHostToDevice);

Note as novas variáveis ponteiro prefixadas com 'd_'. Fiz isto para identificar variáveis que irão ficar no 'device', ou seja, na GPU. Isto não é necessário, está apenas para deixar o código mais claro.
É alocado espaço para estas variáveis na GPU através do comando 'cudaMalloc', que funciona muito similar ao comando 'malloc'.
O método 'cudaMemcpy' possui a seguinte assinatura (neste exemplo):
void cudaMemcpy(endereco destino, endereco fonte, tamanho dados, tipo de transferencia)

• Endereço destino: como queremos copiar esse endereço para a GPU, usamos o nosso ponteiro novo prefixado com 'd_'
• Endereço fonte: onde estão os dados do vetor
• Tamanho dados: quantidade de dados a serem copiados
• Tipo de transferencia: indica a direção que os dados serão transferidos, pode ser (cudaMemcpyHostToDevice, cudaMemcpyDeviceToDevice,cudaMemcpyDeviceToHost,cudaMemcpyHostToHost)

Precisamos agora chamar o nosso método de adição de vetores na GPU, fazemos então:

adicionar_vetor<<<512,>>>(d_a,d_b,d_c,n);

Note quais ponteiros foram utilizados, isto é muito importante. Os dois números 512,512 significam o número de blocos e threads por bloco, respectivamente. Escolher estes números é algo relativo ao problema, apenas escolhi 512 e 512 porque este código poderá somar um vetor de até 512*512=262144 elementos.

Está quase pronto, apenas necessitamos pegar os resultados de volta para a CPU para efetuar uma comparação:

//Obtem os resultados de volta
cudaMemcpy(resultados_gpu, d_c, sizeof(float) * n, cudaMemcpyDeviceToHost);

//verifica se o resultado está ok
for (int j = 0; j < esperado =" a[j]" i="%i">/lib" -I"/common/inc" adicao_vetores.cu -o adicao_vetores