Algorithme de tri utilisant CUDA

(ashwanirathee.com)

3 points par GN⁺ 2025-03-13 | 2 commentaires | Partager sur WhatsApp

Comment améliorer les performances d’un algorithme de tri via le calcul parallèle avec CUDA
Exploration du potentiel de gain de performances en implémentant un tri par fusion (merge sort) de base avec CUDA

Tri par fusion récursif de base (implémentation CPU)

Algorithme de tri qui divise le tableau en deux sous-tableaux, les trie séparément, puis les fusionne
Le tableau est divisé récursivement, puis l’opération de fusion est effectuée lorsque la taille atteint 1
Points clés de l’implémentation
- Utilisation de uint8_t → réduit l’usage mémoire avec de petites valeurs (0 à 255)
- Utilisation de long long → permet de traiter de grands tableaux (jusqu’à 10¹⁸)
- Validation des résultats avec std::sort pour comparer les performances
- Complexité temporelle : O(n log n) en moyenne
- Complexité spatiale : O(n)

Tri par fusion récursif de base avec CUDA

Suit le même schéma que l’implémentation CPU
Implémentation de l’opération de fusion pour une exécution parallèle avec CUDA
Points clés de l’implémentation
- Utilisation de cudaMalloc, cudaMemcpy, cudaFree → allocation de mémoire GPU et transfert de données
- merge<<<1, 1>>>(...) → exécution parallèle de l’opération de fusion
- cudaDeviceSynchronize() → synchronisation jusqu’à la fin de la fusion
- Problème de performance → CUDA étant inefficace pour la récursion, une approche itérative est nécessaire

Comparaison entre les implémentations CPU et GPU

Baisse de performances car les appels récursifs s’exécutent sur le CPU
Avec CUDA, les appels récursifs posent des problèmes de taille de pile et de surcharge d’exécution des kernels
Méthode d’amélioration des performances : passer à une approche itérative (bottom-up)

Tri par fusion itératif bottom-up (implémentation CPU)

Fusion progressive à partir de petits sous-tableaux → plus efficace avec CUDA
La taille des blocs à fusionner augmente selon 1, 2, 4, 8, …

Structure principale du code

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

Points clés de l’implémentation
- Si la taille du tableau n’est pas un multiple de 2, le problème est résolu en limitant les index
- Les opérations de fusion sont exécutées dans une boucle
- Fort potentiel d’amélioration des performances

Tri par fusion itératif bottom-up (implémentation CUDA)

Amélioration des performances en exécutant le tri par fusion itératif en parallèle
Exécution après calcul du nombre de threads et de blocs pour paralléliser l’opération de fusion

Structure principale du code

  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));  
  }

Points clés
- flipflop → alterne l’emplacement de stockage du résultat de la fusion
- gridSize, THREADS_PER_BLOCK → parallélisation des opérations de fusion
- mergeKernel → assigne à chaque thread une opération de fusion propre
- Gestion des index via le calcul des index des threads et des blocs

Résultats de performance

Tri par fusion bottom-up (CUDA) → amélioration nette des performances
- Petits tableaux → le CPU est plus rapide
- Grands tableaux → CUDA prend l’avantage en performances
thrust::sort → meilleures performances GPU sur les grands tableaux
Le gain de performances de CUDA est limité par la surcharge des transferts de données

Conclusion et travaux futurs

Amélioration réussie des performances du tri par fusion basé sur CUDA
Principaux enseignements :
- Compréhension des concepts de traitement parallèle de CUDA et des stratégies de réglage des performances
- Le tri par fusion itératif est plus efficace avec CUDA qu’une approche récursive
- Identification de goulets d’étranglement propres à CUDA, comme la synchronisation des threads et les transferts mémoire
Améliorations futures :
- Séparer et optimiser davantage le travail entre CPU et GPU
- Tester les performances sur des tableaux plus grands
- Combiner thrust::sort avec une implémentation personnalisée
- Optimiser les performances via l’utilisation de mémoire partagée

2 commentaires

kandk 2025-03-14

Il semble que cela ait été implémenté en radix sort dans CUDA, et j’ai déjà eu l’expérience de reproduire exactement la même implémentation en m’en inspirant.

GN⁺ 2025-03-13

Avis Hacker News

Il ne s’agit pas d’une méthode rapide pour trier sur GPU. L’algorithme le plus rapide sur CUDA est Onesweep, qui utilise des techniques complexes pour exploiter le parallélisme du GPU et dépasser ses limites
- Linebender travaille actuellement à appliquer ces idées de manière plus portable sur GPU
- Des ressources associées sont disponibles sur la page wiki de Linebender
Comme d’autres l’ont signalé, cet algorithme n’est pas approprié. Des algorithmes comme Onesweep sont impressionnants, mais difficiles à comprendre
- En regardant le tri par base, qui est l’algorithme central, c’est plus facile à saisir
- Le tri par base peut être implémenté très simplement en parallèle, et constitue une approche belle et élégante
C’est un problème-jouet intéressant. En utilisant les options de coordination des threads, il pourrait y avoir des gains de performance
- Il est aussi amusant d’utiliser Nsight pour identifier les facteurs qui influencent les performances
- Il faut également envisager d’autres algorithmes de tri. Les réseaux de tri comme le tri bitonique demandent davantage de travail, mais peuvent s’exécuter plus vite sur du matériel parallèle
- Une implémentation simple a permis de trier 10M en environ 10 ms sur un H100. Avec davantage de travail, cela pourrait aller plus vite
Le langage Futhark permet d’utiliser plus commodément ce type d’algorithmes sur GPU
- C’est un langage de très haut niveau qui se compile en instructions GPU et peut être utilisé via une bibliothèque Python
- Le site web propose un exemple d’implémentation de tri par fusion
Cela rappelle un petit projet universitaire où un tri bitonique avait été implémenté en CUDA
- L’implémentation du tri bitonique est disponible sur GitHub
Les notes expliquant le concept d’indexation des threads GPU sont bonnes
- Elles présentent aussi un billet de blog personnel sur les avantages en performance du tri vectorisé
J’aime cette implémentation rapide de tri par base
- Elle fonctionne facilement avec l’API du pilote Cuda et, contrairement à CUB, ne se limite pas à l’API runtime
- Onesweep y est également inclus, mais je n’ai pas réussi à le faire fonctionner
J’ai essayé de l’utiliser avec Unity, mais je n’ai pas réussi à surmonter le goulot d’étranglement des transferts de données
- Il y a aussi un surcoût lors de l’utilisation de compute shaders, mais il n’est pas très important
Pour que le tri sur GPU en vaille la peine, il faut de grands tableaux
- Le transfert des données entre la RAM et le GPU prend du temps
Pour résumer et faire gagner du temps : quelqu’un a écrit un algorithme de tri pour GPU, mais il était lent
- Ce n’est pas l’état de l’art, et on ne sait pas clairement si l’auteur maîtrise vraiment l’utilisation efficace du GPU
- Ce n’est qu’une expérimentation personnelle en programmation GPU