CUDA - Meilleure occupation vs moins d'accès mémoire mondial?

Question

Mon code CUDA doit fonctionner avec (réduire en moyenne/std, calculer l'histogramme) 4 tableaux, chaque 2048 flottants de long et déjà stockés dans la mémoire de l'appareil à partir des noyaux précédents.

Il est généralement conseillé de lancer au moins autant de blocs que j'ai de multiprocesseurs. Dans ce cas, cependant, je peux charger chacune de ces matrices dans la mémoire partagée d'un seul bloc et ne lancer que 4 blocs. Ceci est loin de «garder le gpu occupé» mais si j'utilise plus de blocs, je devrai faire plus de communications inter-blocs via la mémoire globale et je prévois que toute utilisation supplémentaire des multiprocesseurs sera en veine en raison de l'extra supplémentaire temps passé à transférer des données dans et hors de la mémoire globale.

Quelle est la meilleure façon de paralléliser dans ce genre de situation?

Answer 1

La communication interbloc n'est pas recommandée dans CUDA. En outre, Fermi prendra en charge l'exécution simultanée du noyau, de sorte qu'une occupation plus élevée deviendra moins importante à l'avenir. Je recommanderais donc de le laisser avec une occupation plus faible pour le moment, sauf si la performance est trop faible.

Answer 2

La quantité de travail que vous faites est relativement faible, donc vous devriez probablement coller avec quatre blocs. Il y a toujours un avantage à garder les données locales au GPU pour les noyaux précédents/suivants.

Fermi va autoriser les noyaux concurrents et c'est exactement ce qui en profitera le plus puisque vous pouvez démarrer le prochain noyau pour occuper les autres SM pendant l'exécution de ce noyau. Cependant, cela suppose qu'il n'y a pas de dépendances entre les deux noyaux - naturellement, vous ne pourrez pas démarrer un nouveau noyau dépendant du résultat du noyau précédent avant la fin du noyau précédent.

Answer 3

Je ne pense pas que vous ayez besoin de tous les 2048 flottants à la fois, et si vous devez réduire, vous pouvez diviser les tableaux dans différentes parties, puis fusionner le résultat à la fin de l'exécution du bloc. Pouvez-vous montrer un exemple de code?

Answer 4

This example montre comment calculer toutes les "statistiques récapitulatives" dans une seule réduction avec Thrust. Un autre exemple montre comment calculer un histogramme en utilisant thrust::sort.