Utiliser la mémoire partagée pour les éléments de tableau voisins?

Question

Je voudrais traiter une image avec CUDA. La nouvelle valeur de chaque pixel est calculée en fonction des deux pixels voisins d'une ligne. Serait-il logique d'utiliser __shared__ mémoire pour les valeurs de pixels, puisque chaque valeur ne sera utilisée que deux fois? Les carreaux ne sont-ils pas aussi la mauvaise façon de le faire, car cela ne convient pas à la structure du problème? Mon approche serait d'exécuter un thread sur chaque pixel et charger les valeurs de pixels voisins chaque fois pour chaque thread.

Answer 1

Toutes les architectures CUDA actuellement supportées ont des caches. A partir de la capacité de calcul 3.5, elles sont particulièrement efficaces pour les données en lecture seule (les données en lecture-écriture étant uniquement mises en cache dans L2, le cache L1 est limité aux données en lecture seule). Si vous marquez le pointeur sur les données d'entrée comme const __restrict__, le compilateur le chargera probablement via the L1 texture cache. Vous pouvez également forcer cela en utilisant explicitement le __ldg() builtin.

Bien qu'il soit possible de gérer explicitement la réutilisation de données provenant de pixels voisins par l'intermédiaire de la mémoire partagée, vous trouverez probablement cela sans aucun bénéfice sur la simple mise en mémoire cache.

Bien sûr, que vous utilisiez ou non la mémoire partagée, vous voulez optimiser la taille des blocs dans la direction x et utiliser un blockSize.y de 1 pour une localisation d'accès optimale.

Answer 2

Combinez en utilisant la mémoire partagée en profitant des accès mémoire coalisés. Tout ce que vous devez faire est de vous assurer que l'image est stockée en ligne. Chaque bloc traiterait un morceau de tableau linéaire. En raison de la réutilisation des données (chaque pixel sauf le premier et le dernier participeraient au traitement trois fois), il serait utile de copier les valeurs de tous les pixels qui seront traités dans la mémoire partagée au début de votre noyau.