Éviter les opérations mov inutiles en ligne PTX

Question

Lors de l'écriture PTX dans un fichier séparé, un paramètre du noyau peut être chargé dans un registre avec:

.reg .u32 test; 
ld.param.u32 test, [test_param]; 

Cependant, lors de l'utilisation PTX en ligne, l'En utilisant Inline Assemblée PTX dans La note d'application CUDA (version 01) décrit une syntaxe dans laquelle le chargement d'un paramètre est étroitement lié à une autre opération. Il fournit cet exemple:

asm("add.s32 %0, %1, %2;" : "=r"(i) : "r"(j), "r"(k)); 

qui génère:

ld.s32 r1, [j]; 
ld.s32 r2, [k]; 
add.s32 r3, r1, r2; 
st.s32 [i], r3; 

Dans de nombreux cas, il est nécessaire de séparer les deux opérations. Par exemple, on peut vouloir stocker le paramètre dans un registre en dehors d'une boucle, puis réutiliser et modifier le registre à l'intérieur d'une boucle. La seule façon que j'ai trouvé pour faire cela est d'utiliser une instruction supplémentaire mov, pour déplacer le paramètre du registre dans lequel il a été implicitement chargé, à un autre registre que je pourrai utiliser plus tard.

Existe-t-il un moyen d'éviter cette instruction mov supplémentaire lors du déplacement de PTX dans un fichier séparé vers PTX intégré?

Answer 1

Si j'étais vous, je ne m'inquiéterais pas trop de ces opérations de mov. N'oubliez pas que PTX n'est pas le code d'assemblage final. PTX est ensuite compilé dans CUBIN avant le lancement du noyau. Entre autres, cette dernière étape effectue l'allocation de registre et supprimera toutes les opérations inutiles mov.

En particulier, si vous passez de %r1 à %r2 et ne jamais utiliser %r1 du tout, l'algorithme est susceptible d'affecter %r1 et %r2 au même registre de matériel et de supprimer le mouvement.