Forcer CUDA pour utilisation vous inscrire pour une variable

J'ai beaucoup de inutilisés registres dans mon noyau. Je voudrais vous dire, CUDA pour utiliser un peu de registres permettant de contenir des données, plutôt que de faire une base de données mondiale de lire à chaque fois que j'en ai besoin. (Je ne suis pas en mesure d'utiliser le mem.)

__global__ void simple(float *gData) {
float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
//work on the data here
}

compiler w/: nvcc -arch sm_20 --ptxas-options=-v simple.cu, et je reçois
0 octets de la trame de pile, 0 octets déversement de magasins, 0 octets déversement des charges
Utilisé 2 registres, 40 octets cmem[0]

__global__ void simple(float *gData) {
register float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
//work on the data here
}

registre déclaration ne fait rien.
0 octets de la trame de pile, 0 octets déversement de magasins, 0 octets déversement des charges
Utilisé 2 registres, 40 octets cmem[0]

__global__ void simple(float *gData) {
volatile float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
//work on the data here
}

volatils déclaration crée de la pile de stockage:
4096 octets de la trame de pile, 0 octets déversement de magasins, 0 octets déversement des charges
21 les résistants, 40 octets cmem[0]

1) Est-il un moyen facile de dire que le compilateur utilise le registre de l'espace pour une variable?
2) Où est la "stack frame': inscrivez-vous, global mem, local mem,...? Qu'est ce qu'un cadre de pile? (Depuis quand les GPU ont une pile? Une pile virtuelle?)
3) Le simple.fichier ptx est essentiellement vide: (nvcc -arch sm_20 -ptx simple.cu)

.loc 2 14 2
ret;

Une idée d'où je peux trouver de la vraie machine/code compilé?

Le compilateur optimisé l'ensemble du code de loin, parce qu'il ne modifie pas la non-état transitoire.
Demander 1024 registres par thread est un joli défi de taille. La plupart des noyaux nécessitent ~des dizaines de registres par thread. Si vous voulez être absolument sûr que le compilateur peut utiliser un registre pour une variable, il doit être un scalaire (c'est à dire, pas un tableau, vous indice dans un for boucle).
La où/quoi frame de pile réponse se trouve ici: stackoverflow.com/questions/7810740/...
flotteur a1,a2,a3,a4,a5; // chaque " un " reçoit un reg . . . . . . . . . . . . volatile float a,b1, b2,b3,b4,b5; // chaque 'b' est sur la pile (local mem) . . . . . . . . . . . Le "volatile" déclaration ne fait rien pour les reg de répartition, mais il crée un local-mem pile

OriginalL'auteur Doug | 2012-08-28

  1. 20
    • Dynamiquement des tableaux indexés ne peuvent pas être stockées dans les registres, parce que le GPU fichier de registre est pas dynamiquement adressable.
    • Variables scalaires sont automatiquement stockées dans les registres par le compilateur.
    • Statiquement indexés (par exemple, lorsque l'indice peut être déterminé au moment de la compilation), petit tableaux (disons, moins de 16 flotteurs) peut être stockées dans les registres par le compilateur.

    SM 2.0 Gpu (Fermi) uniquement en charge jusqu'à 63 registres par thread. Si ce nombre est dépassé, les valeurs d'un registre sera déversé/rempli de locaux (hors de la puce de mémoire, soutenu par la hiérarchie du cache. SM 3.5 Gpu étendre c'jusqu'à 255 registres par thread.

    En général, comme Jared mentionne, en utilisant trop de registres par thread n'est pas souhaitable, car elle réduit l'occupation, et donc réduit le temps de latence cacher capacité dans le noyau. Gpu prospérer sur le parallélisme et le faire en couvrant la latence de la mémoire avec le travail des autres threads.

    Par conséquent, vous ne devriez probablement pas d'optimiser les tableaux dans des registres. Au lieu de cela, vous assurer que votre accès à la mémoire de ces tableaux dans les threads sont aussi proches séquentielle que possible afin de maximiser la coalescence (c'est à dire minimiser la mémoire des transactions).

    L'exemple que vous donnez peut être une affaire de mémoire partagée si:

    1. Nombre de threads dans le bloc utilisent les mêmes données, ou
    2. La par thread taille de la matrice est assez petit pour allouer suffisamment d'espace pour tous les threads dans plusieurs thread se bloque (1024 flotteurs par thread est de loin beaucoup).

    Comme njuffa mentionné, la raison de votre noyau utilise seulement 2 registres est parce que vous ne faites pas quelque chose d'utile avec les données dans le noyau, et le code mort a été tous éliminés par le compilateur.

    Vous êtes ce qui suggère qu'il existe une limite au nombre de reg est qu'un thread peut utiliser (63 pour SM_20). D'où vient-il? Les propriétés de l'appareil montre une limite au nombre de reg par BLOC (regsPerbBock).
    Il s'agit de l'architecture, et le compilateur se charge de veiller à ce pas de numéro de registre supérieure à la limite est utilisé dans le code binaire généré. Un utilisateur n'a pas besoin de s'inquiéter à propos de cette limite autre que pour des raisons de performances (pour comprendre la cause de registre de renverser, par exemple), c'est pourquoi il n'est pas nécessaire de l'inscrire dans la deviceProps structure.
    À l'aide de nombreux registres peuvent être souhaitable parce que la maximisation de l'occupation n'est pas le seul moyen de masquer la latence. Un autre moyen de masquer la latence est l'instruction au niveau de parallélisme. C'est parfois le seul moyen d'atteindre des performances de pointe. Vérifier Vassili Volkov diaporama où l'auteur a obtenu des performances de pointe à seulement 8% d'occupation.

    OriginalL'auteur harrism

  2. 3

    Comme déjà mentionné, les registres (et le PTX "param espace") ne peuvent pas être indexées de façon dynamique. Dans ce but que le compilateur aurait à émettre de code pour un switch...case bloc de transformer la dynamique de l'index dans un immédiat. Je ne suis pas sûr qu'il a déjà fait automatiquement. Vous pouvez aider à y arriver en utilisant une taille fixe n-uplet de la structure et un switch...case. C/C++ métaprogrammation est susceptible d'être l'arme de choix pour garder le code comme cela gérable.

    Aussi, pour CUDA 4.0 utiliser le commutateur de ligne de commande -Xopencc=-O3 afin d'avoir quelque chose, mais la plaine des scalaires (tels que les structures de données) mappés à des registres (voir ce post). Pour CUDA > 4.0, vous devez désactiver le débogage (sans -G option de ligne de commande - optimisation se produit uniquement lorsque le débogage est désactivé).

    PTX niveau permet beaucoup plus de virtuel les registres que le matériel. Ceux qui sont mappés à des registres du matériel au moment du chargement. Le registre limite vous spécifiez vous permet de définir une limite supérieure sur le matériel, les ressources utilisées par le binaire généré. Il sert comme une heuristique pour le compilateur, afin de décider quand un déversement (voir ci-dessous) les registres lors de la compilation de PTX déjà si certains besoins de concurrence peut être atteint (voir "lancement des limites", "occupation" et "concurrent de noyau d'exécution" dans le CUDA Documentation, vous pouvez également profiter de cette présentation des plus intéressantes).

    De Fermi Gpu il y a au plus 64 registres du matériel. La 64e (ou le dernier - lors de l'utilisation de moins que le matériel est maximale) est utilisé par l'ABI que le pointeur de pile, et ainsi de registre "renversement" (c'est à dire en libérant des registres en stockant temporairement leurs valeurs sur la pile et qui se produit quand plusieurs registres sont nécessaires que ce qui est disponible), il est donc intouchable.

    Le lien sur -Xopencc=-O3 est parti et je ne peux pas vraiment trouver des références que dans le contexte de CUDA. Pourriez-vous m'indiquer à certaines ressources ou de l'expliquer, si le comportement avec les récentes cuda (7.0/7.5) est similaire?

    OriginalL'auteur Dude