Réduction du nombre de registres utilisés dans le noyau Cuda

C'est vraiment difficile à dire, le compilateur NVCC n'est pas très intelligent à mon avis.
Vous pouvez essayer des choses évidentes, par exemple en utilisant courtes au lieu d'INT, de passer et d'utiliser des variables par référence (par exemple et variable), des boucles déroulantes, à l'aide de modèles (comme dans C ++). Si vous avez des divisions, des fonctions transcendantaires, ont été appliquées en séquence, essayez de les faire en boucle. Essayez de vous débarrasser des conditionnels, éventuellement les remplacer avec des calculs redondants.

Si vous publiez du code, vous obtiendrez peut-être des réponses spécifiques.

Le nombre d'instructions augmente en cas d'abaissement de l'utilisation de registres d'une explication simple. Le compilateur pourrait utiliser des registres pour stocker les résultats de certaines opérations qui sont utilisées plus d'une fois par votre code afin d'éviter de recalculer ces valeurs, lorsqu'il est forcé d'utiliser moins de registres, le compilateur décide de recalculer ces valeurs qui seraient stockées dans des registres sinon.

Ce n'est généralement pas une bonne approche pour minimiser la pression du registre. Le compilateur optimise un bon travail optimisant les performances globales du noyau projetées et prend en compte beaucoup de facteurs, incluant le registre.

Comment fonctionne-t-il lorsque la réduction des registres a causé une vitesse plus lente

Très probablement Le compilateur a dû renverser des données d'enregistrement insuffisantes dans la mémoire "locale", qui est essentiellement la même que la mémoire globale, et donc très lente

À des fins d'optimisation, je recommanderais d'utiliser des mots-clés tels que const, volatile et ainsi de suite, le cas échéant, pour aider le compilateur à la phase d'optimisation.

Quoi qu'il en soit, ce n'est pas ces minuscules problèmes tels que des registres qui rendent souvent lentement les noyaux Cuda. Je recommanderais d'optimiser le travail avec la mémoire globale, le modèle d'accès, la mise en cache dans la mémoire de texture si possible, des transactions sur la PCIe.

Utilisation de la mémoire partagée que le cache peut mener moins d'un enregistrement d'utilisation et d'empêcher son registre de renversement à la mémoire locale ...

pense que le noyau calcule certaines valeurs et ces valeurs calculées sont utilisées par tous les threads, P> XXX PRE>

Ainsi, au lieu de garder Reg et Reg0 comme des registres et de les rendre possible à la mémoire locale (Mémoire globale), nous pouvons utiliser la mémoire partagée. P>

__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;
}