Introdução aos Streams CUDA para Execução Concorrente
Streams CUDA são uma ferramenta essencial para alcançar concorrência e sobreposição de operações em aplicações de computação acelerada por GPU. Eles permitem que múltiplas operações de kernel, transferências de dados (Host-para-Device e Device-para-Host), e outras tarefas CUDA sejam enfileiradas e executadas de forma assíncrona, potencialmente em paralelo. Este mecanismo é análogo ao cnoceito de pipelines de CPU, onde diferentes estágios de processamento podem operar simultaneamente, aumentando significativamnete a vazão e a eficiência da aplicação.
Considere um cenário onde a mesma função de kernel precisa ser executada várias vezes com conjuntos de dados diferentes. Em vez de executar essas chamadas sequencialmente, o uso de streams CUDA permite que elas sejam despachadas para a GPU quase simultaneamente, com a GPU orquestrando a execução quando recursos estão disponíveis, otimizando o uso do hardawre.
Definição do Kernel de Processamento
Abaixo, apresentamos uma função de kernel CUDA simples que será utilizada para demonstrar a execução concorrente. Este kernel calcula a média ponderada de um elemento e seus dois vizinhos próximos em dois arrays de entrada, e armazena o resultado em um array de saída. A lógica de vizinhança é definida por uma janela circular de tamanho fixo, independentemente da posição do índice principal.
#define TAMANHO_JANELA_VIZINHOS 256 // Define a janela para o cálculo de vizinhos próximos
__global__ void processarDadosKernel(int* dadosEntradaA, int* dadosEntradaB, int* dadosSaida, int numElementos) {
int idxGlobal = threadIdx.x + blockIdx.x * blockDim.x;
if (idxGlobal < numElementos) {
// Calcula índices de vizinhos dentro de uma janela circular fixa
int idxVizinho1 = (idxGlobal + 1) % TAMANHO_JANELA_VIZINHOS;
int idxVizinho2 = (idxGlobal + 2) % TAMANHO_JANELA_VIZINHOS;
// Calcula a média dos três elementos para cada entrada (elemento atual + 2 vizinhos fixos)
float mediaA = (static_cast<float>(dadosEntradaA[idxGlobal]) + dadosEntradaA[idxVizinho1] + dadosEntradaA[idxVizinho2]) / 3.0f;
float mediaB = (static_cast<float>(dadosEntradaB[idxGlobal]) + dadosEntradaB[idxVizinho1] + dadosEntradaB[idxVizinho2]) / 3.0f;
// Calcula a média das duas médias resultantes e armazena como inteiro
dadosSaida[idxGlobal] = static_cast<int>((mediaA + mediaB) / 2.0f);
}
}
Implementação com Streams CUDA
O exemplo a seguir demonstra como utilizar dois streams CUDA para processar grandes volumes de dados em blocos, de forma concorrente. O processo envolve a alocação de memória no host (CPU) e no device (GPU), a inicialização dos dados de entrada, a transferência assíncrona de dados para a GPU, o lançamento assíncrono do kernel e a cópia assíncrona dos resultados de volta para o host.
#include <iostream>
#include <cstdlib> // Para rand()
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define ELEMENTOS_POR_CHUNK (1024 * 1024) // Número de elementos processados por chunk em cada stream
#define MULTIPLICADOR_DADOS_TOTAIS 20 // Fator para definir o tamanho total dos dados
#define TOTAL_DADOS_PROCESSAR (ELEMENTOS_POR_CHUNK * MULTIPLICADOR_DADOS_TOTAIS) // Tamanho total de dados
#define TAMANHO_JANELA_VIZINHOS 256 // Deve ser consistente com a macro do kernel
// Declaração do kernel para o compilador do host
__global__ void processarDadosKernel(int* dadosEntradaA, int* dadosEntradaB, int* dadosSaida, int numElementos);
int main() {
// Declaração das handles dos streams CUDA
cudaStream_t streamProcessamento0, streamProcessamento1;
// Criação dos streams para orquestrar operações assíncronas
cudaStreamCreate(&streamProcessamento0);
cudaStreamCreate(&streamProcessamento1);
// Ponteiros para memória do dispositivo (GPU) para dois conjuntos de buffers de trabalho
int *d_bufferA_s0, *d_bufferB_s0, *d_bufferC_s0;
int *d_bufferA_s1, *d_bufferB_s1, *d_bufferC_s1;
// Alocação de memória na GPU para o Stream 0
cudaMalloc((void**)&d_bufferA_s0, ELEMENTOS_POR_CHUNK * sizeof(int));
cudaMalloc((void**)&d_bufferB_s0, ELEMENTOS_POR_CHUNK * sizeof(int));
cudaMalloc((void**)&d_bufferC_s0, ELEMENTOS_POR_CHUNK * sizeof(int));
// Alocação de memória na GPU para o Stream 1
cudaMalloc((void**)&d_bufferA_s1, ELEMENTOS_POR_CHUNK * sizeof(int));
cudaMalloc((void**)&d_bufferB_s1, ELEMENTOS_POR_CHUNK * sizeof(int));
cudaMalloc((void**)&d_bufferC_s1, ELEMENTOS_POR_CHUNK * sizeof(int));
// Ponteiros para memória do host (CPU) fixada (pinned memory)
// A memória fixada permite transferências de dados assíncronas e mais eficientes entre host e device.
int *h_dadosOrigemA, *h_dadosOrigemB, *h_resultadosFinaisC;
cudaHostAlloc((void**)&h_dadosOrigemA, TOTAL_DADOS_PROCESSAR * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**)&h_dadosOrigemB, TOTAL_DADOS_PROCESSAR * sizeof(int), cudaHostAllocDefault);
cudaHostAlloc((void**)&h_resultadosFinaisC, TOTAL_DADOS_PROCESSAR * sizeof(int), cudaHostAllocDefault);
// Inicialização dos dados de entrada no host com valores aleatórios
for (int i = 0; i < TOTAL_DADOS_PROCESSAR; ++i) {
h_dadosOrigemA[i] = rand() % 1000;
h_dadosOrigemB[i] = rand() % 1000;
}
// Definição das dimensões da grade e do bloco para o lançamento do kernel
dim3 gridDim(ELEMENTOS_POR_CHUNK / TAMANHO_JANELA_VIZINHOS); // Calcula o número de blocos
dim3 blockDim(TAMANHO_JANELA_VIZINHOS); // Define threads por bloco (256, correspondendo à janela)
// Loop principal para processar os dados totais em segmentos duplos usando dois streams
// Em cada iteração, processamos 2 * ELEMENTOS_POR_CHUNK elementos.
for (int i = 0; i < TOTAL_DADOS_PROCESSAR; i += (ELEMENTOS_POR_CHUNK * 2)) {
// --- Operações para o Stream 0 ---
// Cópia assíncrona do primeiro chunk de dados Host -> Device
cudaMemcpyAsync(d_bufferA_s0, h_dadosOrigemA + i, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyHostToDevice, streamProcessamento0);
cudaMemcpyAsync(d_bufferB_s0, h_dadosOrigemB + i, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyHostToDevice, streamProcessamento0);
// Lançamento assíncrono do kernel no Stream 0
processarDadosKernel<<<gridDim, blockDim, 0, streamProcessamento0>>>(d_bufferA_s0, d_bufferB_s0, d_bufferC_s0, ELEMENTOS_POR_CHUNK);
// Cópia assíncrona dos resultados do Device -> Host para o Stream 0
cudaMemcpyAsync(h_resultadosFinaisC + i, d_bufferC_s0, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyDeviceToHost, streamProcessamento0);
// --- Operações para o Stream 1 ---
// Cópia assíncrona do segundo chunk de dados Host -> Device
cudaMemcpyAsync(d_bufferA_s1, h_dadosOrigemA + i + ELEMENTOS_POR_CHUNK, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyHostToDevice, streamProcessamento1);
cudaMemcpyAsync(d_bufferB_s1, h_dadosOrigemB + i + ELEMENTOS_POR_CHUNK, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyHostToDevice, streamProcessamento1);
// Lançamento assíncrono do kernel no Stream 1
processarDadosKernel<<<gridDim, blockDim, 0, streamProcessamento1>>>(d_bufferA_s1, d_bufferB_s1, d_bufferC_s1, ELEMENTOS_POR_CHUNK);
// Cópia assíncrona dos resultados do Device -> Host para o Stream 1
cudaMemcpyAsync(h_resultadosFinaisC + i + ELEMENTOS_POR_CHUNK, d_bufferC_s1, ELEMENTOS_POR_CHUNK * sizeof(int), cudaMemcpyDeviceToHost, streamProcessamento1);
}
// Sincronização dos streams: Bloqueia o host até que todas as operações enfileiradas em cada stream sejam concluídas.
cudaStreamSynchronize(streamProcessamento0);
cudaStreamSynchronize(streamProcessamento1);
// Liberação de todos os recursos de memória alocados e destruição dos streams
cudaFreeHost(h_dadosOrigemA);
cudaFreeHost(h_dadosOrigemB);
cudaFreeHost(h_resultadosFinaisC);
cudaFree(d_bufferA_s0);
cudaFree(d_bufferB_s0);
cudaFree(d_bufferC_s0);
cudaFree(d_bufferA_s1);
cudaFree(d_bufferB_s1);
cudaFree(d_bufferC_s1);
cudaStreamDestroy(streamProcessamento0);
cudaStreamDestroy(streamProcessamento1);
return 0;
}