Optimisation d'un pipeline CUDA de traitement d'images : des bugs aux benchmarks

  • PublicPublic
  • 02 Sep, 2026
Optimisation d'un pipeline CUDA de traitement d'images : des bugs aux benchmarks

Dans cet exemple pratique CUDA, un flux d'images rouges, vertes et bleues est traité : les données sont d'abord copiées du CPU vers le GPU, converties du RGB en niveaux de gris, puis chaque tuile de 32 par 32 pixels est triée afin de pouvoir calculer sa médiane, et enfin les médianes sont copiées vers l'hôte. L'implémentation initiale définit deux kernels — computeRGBToGray pour la conversion des couleurs et computeMedian<TILE_WIDTH, HISTO_SIZE> pour la médiane de chaque tuile, — et main alloue la mémoire, lance les kernels et libère les ressources dans une boucle parallèle OpenMP sur trois images.

Le kernel de médiane cache un défaut subtil mais sérieux. L'écriture en mémoire partagée tile[index] = d_image_gray[index] utilise par erreur un index global pour écrire en mémoire partagée, qui est limitée à un seul bloc de threads. L'exécution du binaire produit une erreur « illegal memory access », et compute-sanitizer indique précisément la cause : une écriture hors limites __shared__ d'un octet à la ligne 55, constatée pour le thread (0,3,0) dans le bloc (20,0,0).

Figure 1. La première partie du pipeline de traitement d'images. D'abord, copier les images RGB de l'hôte vers le device, puis les convertir du RGB en niveaux de gris
Figure 1. La première partie du pipeline de traitement d'images. D'abord, copier les images RGB de l'hôte vers le device, puis les convertir du RGB en niveaux de gris

Pour éviter de telles erreurs d'indexation, NVIDIA a introduit une nouvelle API de lancement dans CCCL. Les kernels sont lancés avec cuda::make_config, cuda::block_dims, cuda::grid_dims et cuda::launch, puis le kernel accepte la configuration comme premier paramètre et utilise cuda::gpu_thread.index(cuda::grid, config) et cuda::gpu_thread.index(cuda::block, config)pour obtenir séparément les indices globaux et de bloc. Les pointeurs bruts peuvent également être remplacés par cuda::std::span, cuda::std::mdspan et cuda::shared_memory_mdspan, qui déclenchent des assertions en mode debug lors d'un accès hors limites du tableau.

Figure 2. La deuxième partie du pipeline de traitement d'images : diviser chaque image en niveaux de gris en tuiles de 32 par 32 pixels, trier les pixels de chaque tuile pour sélectionner sa médiane, puis copier les valeurs médianes depuis d
Figure 2. La deuxième partie du pipeline de traitement d'images : diviser chaque image en niveaux de gris en tuiles de 32 par 32 pixels, trier les pixels de chaque tuile pour sélectionner sa médiane, puis copier les valeurs médianes depuis d

Une fois le code exempt d'erreurs, il est mesuré avec Nsight Systems. L'instrumentation du code avec des plages NVTX — telles que nvtx3::scoped_range et nvtxRangePushA/nvtxRangePop — rend la timeline plus lisible et montre que le kernel de calcul de la médiane domine le profil : il occupe environ 2,1 secondes par image et représente 98,5 % du temps GPU, tandis que les opérations mémoire n'en représentent que 1,5 %.

Figure 3. La timeline Nsight Systems initiale avec annotations NVTX. Les kernels de calcul de la médiane dominent le profil, représentant la quasi-totalité de l'activité GPU et la majeure partie des 6,8 secondes de traitement d'images
Figure 3. La timeline Nsight Systems initiale avec annotations NVTX. Les kernels de calcul de la médiane dominent le profil, représentant la quasi-totalité de l'activité GPU et la majeure partie des 6,8 secondes de traitement d'images

Côté CPU, toute la phase de calcul sur l'image prend 6,8 secondes, et la quasi-totalité de ce temps est consacrée aux kernels de médiane. Cette timeline donne au développeur une cible claire pour les prochaines étapes d'optimisation : le calcul de la médiane est de loin la partie la plus coûteuse du pipeline, et les étapes suivantes du tutoriel se concentrent précisément sur son amélioration.