Algoritmo de ordenação com CUDA
(ashwanirathee.com)- Como melhorar o desempenho de algoritmos de ordenação com computação paralela usando CUDA
- Explora o potencial de ganho de desempenho ao implementar o merge sort básico com CUDA
Merge sort recursivo básico (implementação em CPU)
- Algoritmo de ordenação que divide o array em dois subarrays, ordena cada um deles e depois faz a mesclagem
- Divide o array recursivamente e realiza a operação de mesclagem quando o tamanho chega a 1
- Principais pontos da implementação
- Uso de
uint8_t→ minimiza o uso de memória com valores pequenos (0~255) - Uso de
long long→ permite processar arrays grandes (até 10¹⁸) - Os resultados são validados com
std::sortpara comparação de desempenho - Complexidade de tempo: O(n log n) em média
- Complexidade de espaço: O(n)
- Uso de
Merge sort recursivo básico em CUDA
- Segue o mesmo padrão da implementação em CPU
- A operação de mesclagem foi implementada para rodar em paralelo no CUDA
- Principais pontos da implementação
- Uso de
cudaMalloc,cudaMemcpy,cudaFree→ alocação de memória na GPU e transferência de dados merge<<<1, 1>>>(...)→ executa a operação de mesclagem em paralelocudaDeviceSynchronize()→ sincroniza até a conclusão da mesclagem- Problema de desempenho → como o CUDA é ineficiente para recursão, é necessário uma abordagem iterativa
- Uso de
Comparação entre implementações em CPU e GPU
- Há perda de desempenho porque as chamadas recursivas são executadas na CPU
- Em CUDA, chamadas recursivas geram problemas de tamanho de pilha e overhead de execução de kernel
- Forma de melhorar o desempenho: migrar para uma abordagem iterativa (bottom-up)
Merge sort iterativo bottom-up (implementação em CPU)
- Faz a mesclagem gradualmente a partir de subarrays pequenos → mais eficiente no CUDA
- Aumenta o tamanho do array de mesclagem em 1, 2, 4, 8, … enquanto executa as mesclagens
- Estrutura principal do código
MERGE_SORT(arr, temp, start, end) FOR sub_size ← 1 TO end STEP 2 × sub_size DO FOR left ← 0 TO end STEP 2 × sub_size DO mid ← MIN(left + sub_size - 1, end) right ← MIN(left + 2 * sub_size - 1, end) MERGE(arr, temp, left, mid, right) ENDFOR ENDFOR END MERGE_SORT - Principais pontos da implementação
- Quando o tamanho do array não é múltiplo de 2, o problema é resolvido limitando os índices
- A operação de mesclagem é feita por meio de loops
- Grande potencial de melhoria de desempenho
Merge sort iterativo bottom-up (implementação em CUDA)
- Melhora o desempenho ao executar o merge sort iterativo em paralelo
- A operação de mesclagem é executada após calcular a quantidade de threads e blocos para o paralelismo
- Estrutura principal do código
void mergeSort(uint8_t* arr, uint8_t* temp, long long n) { bool flipflop = true; long long size; for (size = 1; size < n; size *= 2) { numThreads = max(n / (2 * size), (long long)1); gridSize = (numThreads + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; mergeKernel<<<gridSize, THREADS_PER_BLOCK>>>(flipflop ? arr : temp, flipflop ? temp : arr, size, n); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize()); flipflop = !flipflop; } if (!flipflop) CUDA_CHECK(cudaMemcpy(arr, temp, n * sizeof(uint8_t), cudaMemcpyDeviceToDevice)); } - Principais pontos
flipflop→ alterna o local onde o resultado da mesclagem é armazenadogridSize,THREADS_PER_BLOCK→ fazem a paralelização da operação de mesclagemmergeKernel→ atribui uma tarefa de mesclagem única a cada thread- Gerenciamento de índices por meio do cálculo dos índices de threads e blocos
Resultados de desempenho
- Merge sort bottom-up (CUDA) → a melhora de desempenho é clara
- Arrays pequenos → a CPU é mais rápida
- Arrays grandes → o CUDA tem vantagem de desempenho
thrust::sort→ desempenho superior da GPU em arrays grandes- O ganho de desempenho com CUDA é limitado pelo overhead de transferência de dados
Conclusão e trabalhos futuros
- Houve sucesso em melhorar o desempenho do merge sort baseado em CUDA
- Principais aprendizados:
- Aprendizado dos conceitos de processamento paralelo do CUDA e de estratégias de ajuste de desempenho
- Merge sort iterativo → mais eficaz no CUDA do que a abordagem recursiva
- Identificação de gargalos de desempenho específicos do CUDA, como sincronização de threads e transferência de memória
- Próximas melhorias:
- Separação e otimização das tarefas entre CPU e GPU
- Testes de desempenho com arrays ainda maiores
- Combinação de
thrust::sortcom código implementado pelo usuário - Otimização de desempenho com uso de memória compartilhada
2 comentários
Parece que isso foi implementado em CUDA com radix sort; eu também já tive a experiência de implementar exatamente da mesma forma como referência.
Comentários do Hacker News
Não é uma forma rápida de ordenar na GPU. O algoritmo mais rápido em CUDA é o Onesweep, que usa técnicas complexas para aproveitar o paralelismo da GPU e superar limitações
Como em outras opiniões, este algoritmo não é adequado. Algoritmos como o Onesweep são legais, mas difíceis de entender
É um problema de brinquedo interessante para explorar. Usar opções de ajuste de threads pode trazer ganhos de desempenho
A linguagem Futhark torna mais conveniente usar esses algoritmos na GPU
Isso me lembrou de um pequeno projeto da época da faculdade em que implementei bitonic sort em CUDA
As notas que explicam o conceito de indexação de threads na GPU são boas
Gosto de implementações rápidas de radix sort
Tentei usar com Unity, mas não consegui superar o gargalo de transferência de dados
Para valer a pena ordenar na GPU, é preciso ter arrays grandes
Para economizar seu tempo, em resumo: alguém escreveu um algoritmo de ordenação para GPU, mas ele era lento