Anteriormente, abordamos a concorrência em nível de kernel, onde um único kernel é executado em paralelo por múltiplas threads na GPU. Agora, focaremos na concorrência em nível de grid, que permite a execução simultânea de múltiplas tarefas ou kernels através do uso de CUDA Streams.

CUDA Streams: Fundamentos

Toda operação na plataforma CUDA é submetida a um stream, de forma implícita ou explícita. Os streams se dividem em duas categorias fundamentais:

• Stream Padrão (Implícito): Também chamado de stream nulo, é utilizado automaticamente quando nenhum stream é especificado.
• Streams Explícitos (Não nulos): Criados e gerenciados manualmente pelo desenvolvedor.

Para sobrepor diferentes operações CUDA, é obrigatório o uso de streams não nulos. As transferências assíncronas e os lançamentos de kernels nesses streams viabilizam os seguintes cenários de concorrência de granularidade grossa:

1. Sobreposição de computação no host com computação no device.
2. Sobreposição de processamento do host com transferências de dados entre host e device.
3. Sobreposição de transferências de dados com execução de kernels no device.
4. Execução concorrente de múltiplos kernels no device.

Sob a ótica do dispositivo, as operações direcionadas ao stream padrão são processadas estritamente na ordem de emissão. Do lado do host, transferências síncronas bloqueiam a CPU, enquanto o lançamento de kernels é assíncrono por natureza, devolvendo o controle à CPU quase imediatamente. Esse comportamento permite a sobreposição direta de cálculos entre host e device. Para que a tranfserência de dados ocorra de forma assíncrona, é imperativo associá-la explicitamente a um stream dedicado.

// Criação de um stream explícito para gerenciamento manual
cudaError_t status = cudaStreamCreate(cudaStream_t* ptrStream);

// Transferência assíncrona de dados (o stream é o quinto argumento)
// Por padrão, o valor é 0 (stream padrão)
cudaMemcpyAsync(void* destino, void* origem, size_t bytes, cudaMemcpyKind direcao, cudaStream_t stream = 0);

// A memória fixa (pinned memory) é um requisito para transferências assíncronas
cudaMallocHost(void** ptr, size_t tamanho);
cudaHostAlloc(void** ptrHost, size_t tamanho, unsigned int flags);

// Lançamento de kernel em um stream específico
meu_kernel<<<dimgrid dimblock="" memoriacompartilhada="" stream="">>>(argumentos);

// Liberação dos recursos do stream
cudaStreamDestroy(cudaStream_t stream);

// Sincronização e verificação de estado
// cudaStreamSynchronize bloqueia o host até a conclusão de todas as tarefas
// cudaStreamQuery verifica o status sem bloquear a execução do host
cudaStreamSynchronize(cudaStream_t stream);
cudaStreamQuery(cudaStream_t stream);
</dimgrid>

Um padrão arquitetural comum para distribuir operações em múltiplos streams é ilustrado abaixo:

for (int idx = 0; idx < totalStreams; idx++) {
    int deslocamento = idx * bytesPorStream;
    cudaMemcpyAsync(&dadosDevice[deslocamento], &dadosHost[deslocamento], bytesPorStream, cudaMemcpyHostToDevice, listaStreams[idx]);
    processaDados<<<grade bloco="" listastreams="">>>(&dadosDevice[deslocamento]);
    cudaMemcpyAsync(&dadosHost[deslocamento], &dadosDevice[deslocamento], bytesPorStream, cudaMemcpyDeviceToHost, listaStreams[idx]);
}

for (int idx = 0; idx < totalStreams; idx++) {
    cudaStreamSynchronize(listaStreams[idx]);
}
</grade>

Ao analisar a linha do tempo de execução com múltiplos streams, nota-se que, embora as operações estejam em filas distintas, as transferências de dados não ocorrem de forma totalmente paralela. Isso se deve ao barramento PCIe ser um recurso de hardware compartilhado, o que força a serialização dessas transferências. O limite de kernels executados concorrentemente varia conforme a arquitetura do hardware: dispositivos Fermi suportam até 16 vias, enquanto a arquitetura Kepler permite até 32 caminhos simultâneos.

CUDA Events: Monitoramento e Sincronização

Eventos em CUDA atuam como marcadores associados a pontos específicos do fluxo de execução dentro de um stream. Eles são empregados principalmente para duas finalidades:

1. Sincronizar a execução de um stream.
2. Acompanhar e medir o progresso do dispositivo.

// Inicialização dos marcadores
cudaEvent_t evInicio, evFim;
cudaEventCreate(&evInicio);
cudaEventCreate(&evFim);

// Marca o início da contagem
cudaEventRecord(evInicio, 0);

// Executa o processamento na GPU
executaCalculo<<<grade bloco="">>>(parametros);

// Marca o fim da execução
cudaEventRecord(evFim, 0);

// Aguarda a finalização do evento de fim
cudaEventSynchronize(evFim);

// Calcula o tempo decorrido entre os dois eventos
float tempoTotal;
cudaEventElapsedTime(&tempoTotal, evInicio, evFim);

// Limpeza dos recursos
cudaEventDestroy(evInicio);
cudaEventDestroy(evFim);
</grade>