En réduisant le Nombre de Registres Utilisés dans CUDA Noyau

J'ai un noyau qui utilise 17 registres, en la réduisant à 16 m'apporterait 100% d'occupation. Ma question est: existe-il des méthodes qui peuvent être utilisées pour réduire le nombre de registres utilisés, à l'exclusion de la réécriture complète de mon algorithmes de manière différente. J'ai toujours supposé que le compilateur est beaucoup plus malin que je suis, donc par exemple, j'ai souvent l'usage de variables supplémentaires pour plus de clarté seul. Je suis mal dans cette réflexion?

Note: je ne sais à propos de l' --max_registers (ou quelle que soit la syntaxe est le drapeau, mais l'utilisation de la mémoire locale serait plus nuisible que 25% de baisse du taux d'occupation (je devrais tester ce)

Curieusement, j'ai juste essayé maxrregcount=16 et il a en fait réduit le nombre de registres j'ai été en utilisant à 15 et pas de la mémoire locale a été utilisée. Mais c'est en fait devenu plus lent! Comment cela fonctionne?
essayez de profil de votre application. Compilateur peut ainsi introduire des manigances.
Le taux d'occupation est supérieur à 15 registres comme je l'avais prédit et tout le reste est le même, sauf le nombre d'instructions augmente avec l'un registre inférieur comte. à partir de 3,9 à 4,3 mètres de
qu'en est nombre de branches/divergence au cours de l'exécution?
Seulement 8 ans de retard à la fête, mais la raison pour laquelle il a obtenu plus lent est probablement parce que le compilateur a commencé à faire des rematerialisation. Il a évité à l'aide de plusieurs registres juste le calcul de la valeur, encore et encore chaque fois que cela était nécessaire.

OriginalL'auteur zenna | 2010-02-17

  1. 8

    D'occupation peut être un peu trompeuse et 100% d'occupation ne doit pas être votre cible principale. Si vous pouvez obtenir complètement fusionné les accès à la mémoire globale, puis sur un haut de gamme de GPU 50% d'occupation sera suffisante pour masquer la latence de la mémoire globale (pour les flotteurs, même les plus bas pour les doubles). Découvrez la Avancé CUDA C présentation de CG l'année dernière pour plus d'informations sur ce sujet.

    Dans votre cas, vous devez mesurer la performance à la fois avec et sans maxrregcount la valeur 16. La latence de la mémoire locale doit être caché à la suite d'avoir suffisamment de threads, en supposant que vous n'avez pas d'accès aléatoire dans les tableaux (ce qui aurait pour résultat non-fusionné les accès).

    De vous répondre question spécifique à propos de la réduction des registres, poste le code pour des réponses plus détaillées! Comprendre comment les compilateurs de travail en général peut aider, mais n'oubliez pas que nvcc est une optimisation du compilateur avec un grand espace de paramètre, afin de minimiser registre de comptage doit être équilibrée avec la performance globale.

    Comment un 50% d'occupation sera suffisant? Pourriez-vous nous expliquer plus en détails? Merci beaucoup.

    OriginalL'auteur Tom

  2. 6

    Il est vraiment difficile à dire, compilateur nvcc n'est pas très intelligent, à mon avis.

    Vous pouvez essayer des choses évidentes, par exemple à l'aide de courts au lieu de int, l'adoption et l'utilisation des variables par référence (par exemple,&variable), dérouler les boucles, à l'aide de modèles (comme en C++). Si vous avez des divisions, des fonctions transcendantes, été appliquées dans l'ordre, essayez de faire en boucle. Essayez de vous débarrasser de conditions, éventuellement en les remplaçant par des calculs redondants.

    Si tu postes du code, peut-être vous obtiendrez des réponses précises.

    Depuis, les registres de 32 bits, et l'int de 32 bits sur le GPU, ne serait pas de type int et ne court pas faire de différence ?

    OriginalL'auteur Anycorn

  3. 4

    En utilisant la mémoire partagée comme cache peut entraîner moins de registre d'utilisation et de prévenir registre de renverser à la mémoire locale...

    Pense que le noyau calcule des valeurs et ces valeurs calculées sont utilisées par tous les threads,

    __global__ void kernel(...) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int id0 = blockDim.x * blockIdx.x;

    int reg = id0 * ...;
    int reg0 = reg * a / x + y;


    ...

    int val =  reg + reg0 + 2 * idx;

    output[idx] = val > 10;
}

    Donc, au lieu de les garder reg et reg0 que des registres et de les rendre éventuellement de déversement à la mémoire locale (de mémoire globale), on peut utiliser la mémoire partagée.

    __global__ void kernel(...) {
    __shared__ int cache[10];

    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if (threadIdx.x == 0) {
      int id0 = blockDim.x * blockIdx.x;

      cache[0] = id0 * ...;
      cache[1] = cache[0] * a / x + y;
    }
    __syncthreads();


    ...

    int val =  cache[0] + cache[1] + 2 * idx;

    output[idx] = val > 10;
}

    Prendre un coup d'oeil à ce papier pour de plus amples informations..

    Chaque bloc a besoin de sa propre zone de cache et le premier fil de chaque bloc doit le remplir. De sorte que chaque bloc est indépendant et ne nécessite aucune synchronisation. __syncthreads après l'instruction si la synchronisation.'s les threads d'un bloc. Cependant, la série de la partie la hausse de cette façon et peut-être pas une bonne solution..
    Déjà threadidx.x=6 ne sera pas calculer quoi que ce soit. Elle permettra d'obtenir le résultat du calcul à partir du cache et le cache sera le résultat du calcul que le point de synchronisation est transmis. N'est-ce pas?
    Entendez-vous les deux dernières lignes? Lecture dans le cache?? Est-il un moyen pour résoudre ce problème, thread_fence etc?
    Pourriez-vous mettre à jour la source de papier lien? Merci.

    OriginalL'auteur phoad

  4. 2

    Il n'est généralement pas une bonne approche pour minimiser le registre de la pression. Le compilateur fait un bon travail d'optimisation de l'ensemble projeté la performance du noyau, et il tient compte de beaucoup de facteurs, incliding registre.

    Comment fonctionne lorsque la réduction des registres causé un ralentissement de la vitesse de

    Probablement le compilateur a le déversement de l'insuffisance des données du registre en "local" de la mémoire, qui est essentiellement la même que la mémoire globale, et donc très lent

    Pour des raisons d'optimisation je vous recommande d'utiliser comme mots-clés const et volatile, et ainsi de suite si nécessaire, pour aider le compilateur sur l'optimisation de la phase.

    De toute façon, il n'est pas de ces petits problèmes comme les registres qui font souvent des noyaux CUDA courir lentement. Je le recommande pour optimiser le travail avec la mémoire globale, le modèle de l'accès, la mise en cache dans la mémoire de texture, si possible, les transactions sur le PCIe.

    OriginalL'auteur Tosha

  5. 1

    L'instruction augmentation lors de l'abaissement du registre d'utilisation ont une explication simple. Le compilateur pourrait être l'utilisation des registres pour stocker les résultats de certaines opérations qui sont utilisés plus d'une fois par le biais de votre code afin d'éviter de recalculer ces valeurs, quand ils sont forcés d'utiliser moins de registres, le compilateur décide de recalculer les valeurs qui devraient être stockées dans les registres autrement.

    OriginalL'auteur Victor M. Alvarez