Essayer de lancer ceci: https://github.com/Celebrandil/CudaSift sur un NVS4200M, qui est sm_21, non sm_35 au besoin. Le seul problème dans le fonctionnement du projet mentionné est le code (cudaSiftD.cu:205):Peut être remplacé par _shfl_xor pour s'exécuter sur sm_21?
for (int i = 1; i < = 16; i * = 2) + = somme __shfl_xor (somme , je);
Existe-t-il un code équivalent possible?
Eh bien, presque toute intrinsèque CUDA peut être remplacé, donc je vais interpréter votre question
Peut
__shfl_xorêtre remplacé bon marché sur les GPU SM_21?
Et la réponse est: Pas vraiment; vous encourrez une pénalité. Votre meilleure option, comme @ commentaire de RobertCrovella suggère est d'utiliser la mémoire partagée:
__syncthreads())Je ne pas orthographier le code de ne pas prendre le plaisir loin pour vous :-)
modifier: Bien que l'exécution de lecture aléatoire est plus complexe, il est encore, au moins sémantiquement, une opération sur les registres; et il ne nécessite pas de synchronisation. L'alternative de mémoire partagée serait donc plus lente.
Je ne considère pas que shuffle est un cycle d'horloge pour deux raisons: 1) il y a 32 shuffles émettables par cycle sur un multiprocesseur [http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html # arithmetic-instructions], et 2) l'opération shuffle est effectuée par le cache qui gère la mémoire partagée. En substance, l'utilisation de shuffle est environ deux fois plus performante que la mémoire partagée - voir [http: //on-demand.gputechconf.com/gtc/2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf] –
@FlorentDUGUET: Edité pour refléter votre commentaire. Votre lien ne fonctionne pas, je pense que vous avez une faute de frappe avec les parenthèses. – einpoklum
Liens brisés dans le commentaire ci-dessus: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions et http://on-demand.gputechconf.com/gtc/ 2013/presentations/S3174-Kepler-Shuffle-Tips-Tricks.pdf –
Si la question est plus sur la façon de remplacer cet extrait de code par un compatible avec sm_21, vous voudrez peut-être avoir un oeil sur CUB, la partie bloc-réduire here. L'un des paramètres du modèle est l'achitecture de votre appareil. La macro __CUDA_ARCH__ peut vous aider à sélectionner l'implémentation la plus appropriée, voir here.
Oui il y a, si vous voulez l'écrire. – talonmies
Pratiquement tout ce que vous pouvez faire avec les opérations de shuffle peut être fait avec des opérations de mémoire partagée, ce qui permet également une communication inter-thread. Je ne suggère pas que la mise en œuvre soit identique, juste qu'il existe un "code équivalent possible" utilisant la mémoire partagée. –
@talonmies comment ce commentaire peut-il aider l'OP? C'est une question non triviale, car je ne considère pas l'intrinsèque de shuffle comme une simple caractéristique de cuda. –