Algoritmo de ordenamiento con CUDA

• Cómo mejorar el rendimiento de algoritmos de ordenamiento mediante cómputo paralelo usando CUDA
• Se explora la posibilidad de mejorar el rendimiento implementando el merge sort básico con CUDA

Merge sort recursivo básico (implementación en CPU)

• Algoritmo de ordenamiento que divide el arreglo en dos subarreglos, ordena cada uno y luego los fusiona
• Divide el arreglo de forma recursiva y, cuando el tamaño llega a 1, realiza la fusión
• Puntos clave de la implementación
  • Uso de uint8_t → minimiza el uso de memoria con valores pequeños (0~255)
  • Uso de long long → permite procesar arreglos grandes (hasta 10¹⁸)
  • Se valida el resultado con std::sort para comparar el rendimiento
  • Complejidad temporal: promedio O(n log n)
  • Complejidad espacial: O(n)

Merge sort recursivo básico en CUDA

• Sigue el mismo patrón que la implementación en CPU
• Se implementa para ejecutar la operación de fusión en paralelo en CUDA
• Puntos clave de la implementación
  • Uso de cudaMalloc, cudaMemcpy, cudaFree → asignación de memoria en GPU y transferencia de datos
  • merge<<<1, 1>>>(...) → ejecuta la operación de fusión en paralelo
  • cudaDeviceSynchronize() → sincroniza hasta que termine la fusión
  • Problema de rendimiento → CUDA es ineficiente para manejar recursión, por lo que se requiere un enfoque iterativo

Comparación entre implementación en CPU y GPU

• Como las llamadas recursivas se ejecutan en la CPU, se produce degradación del rendimiento
• En CUDA, las llamadas recursivas generan problemas de tamaño de pila y sobrecarga de ejecución del kernel
• Método para mejorar el rendimiento: cambiar a un enfoque iterativo (bottom-up)

Merge sort iterativo bottom-up (implementación en CPU)

• Fusiona progresivamente desde subarreglos pequeños → más eficiente en CUDA
• Fusiona aumentando el tamaño del arreglo de fusión en 1, 2, 4, 8, …
• Estructura principal del 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  
  

• Puntos clave de la implementación
  • Si el tamaño del arreglo no es múltiplo de 2, el problema se resuelve limitando los índices
  • La operación de fusión se realiza mediante bucles
  • Hay gran potencial de mejora de rendimiento

Merge sort iterativo bottom-up (implementación en CUDA)

• Mejora el rendimiento ejecutando en paralelo el merge sort iterativo
• Para realizar la fusión en paralelo, se calculan y lanzan la cantidad de hilos y bloques
• Estructura principal del 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));  
    }  
  

• Puntos principales
  • flipflop → alterna la ubicación donde se guarda el resultado de la fusión
  • gridSize, THREADS_PER_BLOCK → realizan la paralelización de la operación de fusión
  • mergeKernel → asigna una tarea de fusión única a cada hilo
  • Gestión de índices mediante el cálculo de índices de hilos y bloques

Resultados de rendimiento

• Merge sort bottom-up (CUDA) → la mejora de rendimiento es clara
  • Arreglos pequeños → la CPU es más rápida
  • Arreglos grandes → CUDA tiene ventaja de rendimiento
• thrust::sort → buen rendimiento de GPU en arreglos grandes
• La mejora de rendimiento de CUDA está limitada por la sobrecarga de transferencia de datos

Conclusión y trabajo futuro

• Se logró mejorar el rendimiento del merge sort basado en CUDA
• Puntos principales aprendidos:
  • Se aprendieron los conceptos de procesamiento paralelo en CUDA y estrategias de ajuste de rendimiento
  • El merge sort iterativo → es más efectivo en CUDA que el enfoque recursivo
  • Se identificaron cuellos de botella propios de CUDA, como sincronización de hilos y transferencia de memoria
• Mejoras futuras:
  • Separación y optimización del trabajo entre CPU y GPU
  • Pruebas de rendimiento con arreglos más grandes
  • Combinar thrust::sort con código implementado por el usuario
  • Optimización del rendimiento mediante uso de memoria compartida

[kandk]

Parece que estaba implementado con radix sort en CUDA, y recuerdo haber tenido la experiencia de implementarlo exactamente igual como referencia.

[GN⁺]

Opiniones de Hacker News

• No es una forma rápida de ordenar en GPU. El algoritmo más rápido en CUDA es Onesweep, que usa técnicas complejas para aprovechar el paralelismo de la GPU y superar sus limitaciones

  • Linebender está trabajando en aplicar estas ideas de una manera más portable a la GPU
  • Se puede consultar material relacionado en la página wiki de Linebender

• Como dicen otras opiniones, este algoritmo no es adecuado. Algoritmos como Onesweep son geniales, pero difíciles de entender

  • Si se revisa el algoritmo clave, radix sort, es más fácil de comprender
  • Radix sort se puede implementar de forma muy sencilla para paralelización, y es un enfoque bello y elegante

• Está bien como problema de juguete interesante. Usar opciones de ajuste de hilos podría mejorar el rendimiento

  • También es interesante usar Nsight para identificar qué factores afectan el rendimiento
  • También hay que considerar otros algoritmos de ordenamiento. Los sorting networks, como bitonic sort, requieren más trabajo, pero pueden ejecutarse más rápido en hardware paralelo
  • Hice una implementación simple que ordena 10M en unos 10 ms en una H100. Con más trabajo, podría hacerse más rápida

• El lenguaje Futhark permite usar este tipo de algoritmos de forma más cómoda en GPU

  • Es un lenguaje de muy alto nivel que compila a instrucciones de GPU y se puede usar desde una biblioteca de Python
  • En el sitio web hay un ejemplo de implementación de merge sort

• Me recordó a un pequeño proyecto que hice en la universidad implementando bitonic sort en CUDA

  • La implementación de bitonic sort se puede ver en GitHub

• Están buenas las notas que explican el concepto de indexación de hilos en GPU

  • Presenta una publicación de blog personal sobre las ventajas de rendimiento del ordenamiento vectorizado

• Me gusta la implementación rápida de radix sort

  • Funciona fácilmente con la API del driver de Cuda y, a diferencia de CUB, no está limitada a la API de runtime
  • También incluye Onesweep, pero no logré hacerlo funcionar

• Intenté usarlo con Unity, pero no pude superar el cuello de botella de transferencia de datos

  • Incluso al usar compute shaders hay sobrecarga, pero no es tan grande

• Para que valga la pena ordenar en GPU, se necesitan arreglos grandes

  • La transferencia de datos entre la RAM y la GPU toma tiempo

• Para ahorrarles tiempo, el resumen es: alguien escribió un algoritmo de ordenamiento para GPU, pero era lento

  • No es estado del arte, y no está claro si el autor sabe cómo usar la GPU de forma efectiva
  • Solo es un experimento personal de programación en GPU