Faible utilisation de GPU dans CUDA

Question

J'ai implémenté un programme qui utilise différents flux CUDA provenant de différents threads CPU. La copie de mémoire est implémentée via cudaMemcpyAsync en utilisant ces flux. Les lancements de noyau utilisent également ces flux. Le programme fait des calculs en double précision (et je soupçonne que c'est le coupable, cependant, cuBlas atteint 75-85% d'utilisation du processeur pour la multiplication des matrices de double s). Il y a également des opérations de réduction, mais elles sont implémentées via if(threadIdx.x < s) avec s diminuant 2 fois dans chaque itération, donc des déformations bloquées devraient être disponibles pour d'autres blocs. L'application est gourmande en GPU et en CPU, elle commence avec un autre travail dès que le précédent est terminé. Donc je m'attends à ce qu'il atteigne 100% de CPU ou de GPU.

Le problème est que mon programme génère 30-40% de la charge GPU (et environ 50% de la charge CPU), si vous faites confiance à GPU-Z 1.9.0. La charge du contrôleur de mémoire est de 9-10%, la charge de l'interface de bus est de 6%. C'est pour le nombre de threads CPU égal au nombre de cœurs CPU. Si je double le nombre de threads CPU, les charges restent à peu près les mêmes (y compris la charge du processeur).

Alors, pourquoi est-ce? Où est le goulot d'étranglement?

J'utilise GeForce GTX 560 Ti, CUDA 8RC, MSVC++ 2013, Windows 10.

Un je suppose que Windows 10 applique une économie d'énergie agressive, même si les températures GPU et CPU sont faibles, la le plan d'alimentation est réglé sur "haute performance" et l'alimentation est de 700W alors que la consommation d'énergie avec max CPU et GPU TDP est d'environ 550W. Une autre supposition est que la vitesse à double précision est 1/12 de la vitesse à simple précision car il y a 1 cœur CUDA à double précision pour 12 cœurs CUDA à simple précision sur ma carte, et GPU-Z à 100% la situation lorsque tous les cœurs à simple précision et à double précision sont utilisés. Cependant, les chiffres ne correspondent pas tout à fait.

Answer 1

Apparemment, la raison était faible taux d'occupation en raison de threads CUDA en utilisant trop de registres par défaut. Pour indiquer au compilateur la limite du nombre de registres par thread, __launch_bounds__ peut être utilisé, comme décrit here. Donc, pour être en mesure de lancer les 1536 threads dans 560 Ti, pour la taille de bloc 256 ce qui suit peut être spécifié:

_global__ void __launch_bounds__(256, 6) MyKernel(...) { ... } 

Après la limitation du nombre de registres par fil de CUDA, l'utilisation de la GPU a porté à 60% pour moi .

Par ailleurs, 5xx cartes de la série sont toujours pris en charge par nSight v5.1 pour Visual Studio. It can be downloaded à partir du archive.

EDIT: les indicateurs suivants ont encore augmenté l'utilisation de GPU à 70% dans une application qui utilise de multiples flux de GPU à partir de plusieurs threads de processeur:

cudaSetDeviceFlags(cudaDeviceScheduleYield | cudaDeviceMapHost | cudaDeviceLmemResizeToMax); 

• cudaDeviceScheduleYield permet d'autres threads exécuter lorsqu'une CPU fil est attente de l'opération GPU, plutôt que de tourner GPU pour le résultat .
• cudaDeviceLmemResizeToMax, comme je l'ai compris, rend le noyau lance eux-mêmes asynchrone et évite la mémoire locale excessive allocations & désallocations.