Uma Introdução a CUDA¶

O CUDA (Compute Unified Device Architecture) é o modelo de programação paralela desenvolvido pela NVIDIA. Ele permite que desenvolvedores usem o poder das GPUs para acelerar aplicações de alto desempenho, executando milhares de threads simultaneamente.

Em um sistema HPC com SLURM, como o Franky ou o SDumont, a execução de programas CUDA segue três etapas principais:

Carregar o módulo CUDA disponível no cluster.
Compilar o código com o compilador nvcc.
Executar o binário com o srun ou sbatch(solicitando uma GPU).

Preparando o ambiente¶

Liste os módulos disponíveis e carregue o CUDA:

module avail cuda
module load cuda/12.8.1

Verifique se o compilador CUDA está ativo:

nvcc --version

Código Base (CPU)¶

Vamos começar com um programa simples em C++ que soma os elementos de dois vetores na CPU.

Crie o arquivo:

nano exemplo-cpu.cpp

Cole o código abaixo:

#include <iostream>
#include <cmath>
#include <chrono> 

// Função que soma os elementos de dois vetores
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1'000'000'000; 
  float *x = new float[N];
  float *y = new float[N];

  // Inicializa os vetores
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Início da medição de tempo
  auto start = std::chrono::high_resolution_clock::now();

  // Executa a soma
  add(N, x, y);

  // Fim da medição
  auto end = std::chrono::high_resolution_clock::now();
  std::chrono::duration<double> elapsed = end - start;

  // Calcula erro e soma total
  float maxError = 0.0f;
  double somaTotal = 0.0;

  for (int i = 0; i < N; i++) {
    maxError = fmax(maxError, fabs(y[i] - 3.0f));
    somaTotal += y[i];
  }

  std::cout << "Erro: " << maxError << std::endl;
  std::cout << "N: " << N << std::endl;
  std::cout << "Soma total: " << somaTotal << std::endl;
  std::cout << "Tempo de execução: " << elapsed.count() << " segundos" << std::endl;

  delete [] x;
  delete [] y;

  return 0;
}

Compile e execute localmente (sem GPU):

g++ -O3 exemplo-cpu.cpp -o ex-cpu

Para executar:

srun --partition=merry_gpu --mem=1G ./ex-cpu

Saída esperada:

Erro: 0
N: 1000000000
Soma total: 3e+09
Tempo de execução: 0.30549 segundos

Migrando para GPU (CUDA)¶

Agora, vamos reescrever o mesmo código para rodar na GPU.

Crie o arquivo:

nano exemplo-gpu.cu

Analise o código abaixo:

#include <iostream>
#include <cmath>
#include <chrono>
#include <cuda_runtime.h>

using namespace std;

// Kernel CUDA (função que será executada na GPU)
__global__
void add(int n, float *x, float *y)
{
  // Calcula o índice global da thread
   // cada thread obtém um índice único 'i'
  int i = blockIdx.x * blockDim.x + threadIdx.x;

  // Garante que a thread só acesse posições válidas do vetor
  if (i < n)

    // Cada thread realiza a soma de UM elemento:
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1'000'000'000;
  size_t size = N * sizeof(float);

  // Cria os vetores na CPU
  float *x = new float[N];
  float *y = new float[N];

  // Inicializa na CPU
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Cria os vetores para a GPU
  float *d_x, *d_y;
  // reserva o espaço para os vetores na GPU
  cudaMalloc(&d_x, size);
  cudaMalloc(&d_y, size);

  // início do tempo (incluindo transferência)
  auto start = std::chrono::high_resolution_clock::now();

  auto dados_start = std::chrono::high_resolution_clock::now();
  // Transferencia dos dados CPU → GPU
  cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
  auto dados_end = std::chrono::high_resolution_clock::now();

  // Configurando o kernel
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;

  auto start_comp = chrono::high_resolution_clock::now();
  // Computação na GPU
  add<<<numBlocks, blockSize>>>(N, d_x, d_y);

  // Semáforo para aguardar o término da computação na GPU 
  cudaDeviceSynchronize();

  auto end_comp = chrono::high_resolution_clock::now();


  auto start_dados = chrono::high_resolution_clock::now();
  // Transferencia de dados GPU → CPU
  cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);

  auto end_dados = chrono::high_resolution_clock::now();

  // fim do tempo
  auto end = std::chrono::high_resolution_clock::now();
  std::chrono::duration<double> total = end - start;

  std::chrono::duration<double> comp = end_comp - start_comp;
  std::chrono::duration<double> dados_crus = dados_end - dados_start;
  std::chrono::duration<double> dados_comp = end_dados - start_dados;


  // Criação dos das variaveis para exibir o output
  float maxError = 0.0f;
  double somaTotal = 0.0;

  // Preparação dos dados para exibir
  for (int i = 0; i < N; i++) {
    maxError = fmax(maxError, fabs(y[i] - 3.0f));
    somaTotal += y[i];
  }

  std::cout << cudaGetLastError() << std::endl;
  std::cout << "Erro: " << maxError << std::endl;
  std::cout << "N: " << N << std::endl;
  std::cout << "Soma total: " << somaTotal << std::endl;

  std::cout << "Tempo de transferência dos dados (CPU->GPU) " << dados_crus.count() << " segundos" << std::endl;
  std::cout << "Tempo de calculos na GPU " << comp.count() << " segundos" << std::endl;
  std::cout << "Tempo de transferência dos dados (CPU<-GPU) " << dados_comp.count() << " segundos" << std::endl;
  std::cout << "Tempo total de tudo mesmo " << total.count() << " segundos" << std::endl;

  // Libera memória
  cudaFree(d_x);
  cudaFree(d_y);
  delete[] x;
  delete[] y;

  return 0;
}

Compile com o nvcc:

nvcc exemplo-gpu.cu -o ex-gpu

Agora, vamos executar o binário pedindo uma GPU via srun:

srun --partition=merry_gpu --gres=gpu:1 --mem=1G ./ex-gpu

Esse --gres=gpu:1 aloca 1 GPU para executar o binário ./ex-gpu na fila do Cluster Franky

Saída esperada:

0
Erro: 0
N: 1000000000
Soma total: 3e+09
Tempo de transferência dos dados (CPU->GPU) 0.663264 segundos
Tempo de calculos na GPU 0.0722021 segundos
Tempo de transferência dos dados (CPU<-GPU) 0.307847 segundos
Tempo total de tudo mesmo 1.04331 segundos

Por que a soma pode não acelerar a aplicação?¶

Por elemento, a expressão y[i] = x[i] + y[i] realiza uma adição, lê dois valores e escreve um resultado. Para float de quatro bytes, isso corresponde a cerca de 12 bytes de tráfego lógico de memória por adição, sem contar efeitos de cache.

Há pouco cálculo para muitos dados movimentados. Para vetores grandes, esse padrão tende a ser limitado pela largura de banda da memória, ou memory-bound. O kernel ainda pode ser rápido na GPU; as cópias entre CPU e GPU podem dominar o intervalo completo medido.

A versão CPU também acessa memória e pode usar caches; a ausência de transferências explícitas para uma GPU não torna esses acessos gratuitos. Neste programa, não criamos múltiplas threads de CPU.

Uma estratégia comum é copiar os dados para a GPU, realizar várias operações e copiar apenas os resultados necessários ao final. Isso permite amortizar o custo das transferências.

Para entender se você entendeu¶

Varie as threads por bloco

Teste blockSize = 128, 256 e 512, mantendo o cálculo de numBlocks por arredondamento para cima. Verifique a correção e registre os tempos. Mais threads por bloco não garantem menor tempo.

Varie o número de blocos

Não basta substituir numBlocks por 32, 128 ou 1024 no kernel original. Como cada thread processa apenas um elemento, isso pode deixar grande parte do vetor sem calcular. Um tempo menor com resultado incorreto não é uma melhoria.

Faça alterações no código para permitir que cada thread processe vários elementos.

Depois teste 32, 128 e 1024 blocos com 256 threads por bloco. A thread começa em seu índice global e avança pelo total de threads do grid, cobrindo outras posições do vetor.