Rédaction d'une spécification pour Nvidia Kepler (binaires CUDA, version en langage sm_30) pour Ghidra

De nombreuses spécifications pour Ghidra ont déjà été écrites pour les langages de processeur ordinaires, mais rien pour les langages graphiques. C'est compréhensible, car il a ses propres spécificités: des prédicats, des constantes par lesquelles les paramètres sont passés, y compris d'autres choses héritées des shaders. De plus, le format utilisé pour stocker le code est souvent propriétaire et doit être inversé par vous-même.



Dans cet article, à l'aide de deux exemples, nous verrons ce qui est quoi.



Le premier programme est le axpy le plus simple (analogue de Hello World pour GPGPU). Le second permet de comprendre la mise en œuvre des conditions et des sauts sur le GPU, puisque tout y est différent.



Toutes les langues Nvidia utilisent peu d'encodage endian, donc copiez immédiatement les octets de l'éditeur hexadécimal dans un bloc-notes (par exemple, Notepad ++) dans l'ordre opposé de 8 pièces (la longueur des instructions ici est constante). Ensuite, grâce à une calculatrice programmeur (celle de Microsoft convient), nous la traduisons en code binaire. Ensuite, nous cherchons les correspondances, composons le masque de l'instruction, puis les opérandes. Pour décoder et rechercher le masque, l'éditeur hexadécimal et le désassembleur cuobjdump ont été utilisés, parfois l'assembleur est requis, comme dans AMDGPU (car le désassembleur n'y est pas disponible, mais c'est un sujet pour un article séparé). Cela fonctionne comme ceci: nous essayons d'inverser séquentiellement tous les bits suspects dans la calculatrice, puis nous obtenons une nouvelle valeur hexadécimale pour les octets, nous les substituons dans un binaire compilé via nvcc ou un assembleur, s'il existe, ce qui n'est pas toujours le cas.Ensuite, nous vérifions à travers cuobjdump.



Je diffuse la source au format (principalement en C, sans avantages et sans POO pour une connexion plus étroite avec le code GPU de la machine), puis disasm + octets à la fois, car c'est plus pratique, ils n'ont tout simplement pas besoin d'être interchangés.



Copiez-le dans axpy.cu et compilez-le via cmd: nvcc axpy.cu --cubin --gpu-architecture sm_30 Démontez le

fichier ELF résultant nommé axpy.cubin au même endroit: cuobjdump axpy.cubin -sass



Exemple 1:

__global__ void axpy(float param_1, float* param_2, float* param_3) {
unsigned int uVar1 = threadIdx.x;
param_2[uVar1] = param_1 * param_3[uVar1];
}

Exemple 2:

__global__ void predicates(float* param_1, float* param_2) {
    unsigned int uVar1 = threadIdx.x + blockIdx.x * blockDim.x;
    if ((uVar1 > 5) & (uVar1 < 10)) param_1[uVar1] = uVar1;
    else param_2[uVar1] = uVar1;
}

Il n'est pas difficile de deviner que les tests ont été conçus à l'origine pour le code machine de sorte que le compilateur n'avait rien à optimiser. Pour tout le reste, vous devrez annuler manuellement les optimisations. Dans les exemples complexes, cela peut ne pas être possible du tout, donc dans de tels cas, vous devrez vous fier au décompilateur et au frontend.



En général, la règle est la suivante: pour tester le frontend, nous prenons tout exemple simple (avec un minimum d'optimisations possibles) premier exemple approprié (erreurs de reproduction). Pour le reste, le code décompilé sera déjà avec des optimisations (ou juste en quelque sorte corrigé par refactoring). Mais pour l'instant, la tâche principale est au moins juste un code correct qui fait la même chose que le code machine. Il s'agit de la "modélisation logicielle". La «modélisation logicielle» elle-même n'implique pas de refactoring, de traduction de C en C ++, de restauration de classes, et plus encore de choses telles que l'identification de modèles.



Nous recherchons maintenant des modèles pour les mnémoniques, les opérandes et les modificateurs.



Pour ce faire, comparez les bits (en binaire) entre les instructions suspectes (ou les chaînes, si elles sont plus pratiques à appeler). Vous pouvez également profiter de ce que les autres utilisateurs publient sur leurs questions stackoverflow telles que «aider à comprendre le code binaire / sass / machine», utiliser des didacticiels (y compris en chinois) et d'autres ressources. Ainsi, le numéro de l'opération principale est stocké dans les bits 58-63, mais il y a des bits supplémentaires 0-4 (ils distinguent les instructions "I2F", "ISETP", "MOV32I"), quelque part au lieu d'eux 0-2 (par négligence, 3- 4 bits dans des instructions vides, ils sont marqués comme "UNK" dans la spécification).



Pour les registres et les nombres constants, vous pouvez expérimenter avec le désassembleur afin de trouver tous les bits qui affectent la sortie de vidage, comme celui présenté sous le spoiler. Tous les champs que j'ai réussi à trouver sont dans la spécification sur Github, fichier CUDA.slaspec, section token.



Ensuite, vous devez trouver des adresses pour les registres, encore une fois, ils sont sur Github. Ceci est nécessaire car au niveau micro, Sleigh considère les registres comme des variables globales dans un espace de type "register_space", mais depuis Puisque leur espace n'est pas marqué comme «inférable» (et certainement pas), alors ils deviennent dans le décompilateur soit des variables locales (le plus souvent avec l'interfixe «Var», mais parfois le préfixe «local» était aussi semblable) ou des paramètres (le « param_ "). SP n'a jamais été utile, il est principalement nécessaire pour s'assurer que le décompilateur fonctionne. Un PC (quelque chose comme IP de x86) est nécessaire pour émuler.



Ensuite, il y a des registres de prédicats, quelque chose comme des drapeaux, mais plus «à usage général» que pour un objectif prédéterminé, comme le débordement, (non) égal à zéro, etc.

Puis un registre de blocage pour simuler un tas d'instructions ISCADD .CC et IMAD.HI, puisque le premier de mon implémentation effectue un comptage pour lui-même et pour le second, afin d'éviter de transférer une partie de la somme vers les 4 octets supérieurs, puisque cela va gâcher la décompilation. Mais alors vous devez verrouiller le registre suivant jusqu'à ce que l'opération IMAD.HI soit terminée. Quelque chose de similaire, c'est-à-dire l'écart entre la documentation officielle et la sortie attendue du décompilateur était déjà dans le module SPU pour le même Ghidra.



Ensuite, il existe des registres spéciaux, qui sont jusqu'à présent implémentés via cpool. À l'avenir, je prévois de les remplacer par les symboles par défaut pour certains espaces "inférables". Ce sont les mêmes threadIdx, blockIdx.



Ensuite, nous lions les variables aux champs dest, par0, par1, par2, res. Ensuite, il y a des sous-tables, et après elles - de quoi il s'agissait - les tables principales (racine) avec les instructions principales.



Ici, il est nécessaire de suivre strictement le format «opérande mnémotechnique», cependant, un soulagement est donné pour les modificateurs, qui, néanmoins, doivent être attachés aux mnémoniques ou à la section avec des opérandes. Aucun autre format n'est autorisé, même le même asm Hexagon DSP devra être adapté à cette syntaxe, ce qui n'est cependant pas très difficile.



La dernière étape consistera à rédiger l'implémentation des instructions dans le langage de microprogrammation Pcode. La seule chose que je voudrais noter dans le premier exemple est les instructions ISCADD .CC et IMAD.HI, où la première prend un pointeur vers les registres et les déréférence comme pointeurs pour 8 octets au lieu de 4. Ceci est fait intentionnellement afin de mieux s'adapter au décompilateur et son comportement, malgré ce qui est écrit dans la documentation Nvidia sur le transfert d'une partie du montant.



Pour le second exemple, il vaut mieux vérifier les paramètres du décompilateur en face de l'inscription «Simplifier la prédication». Le fait est que les prédicats sont une seule et même condition pour différentes instructions, essentiellement rien de plus que le "SIMD" bien connu, ou son équivalent suivant. Ceux. si le bit de prédicat est positionné, les instructions sont en outre exécutées en ligne.



Vous devez également prendre l'habitude d'écrire immédiatement une implémentation pour chaque instruction, et pas seulement un prototype (mnémoniques d'opérande), car il existe également un décompilateur, un émulateur et d'autres analyseurs.

Mais en général, écrire une implémentation dans Pcode est une tâche encore plus facile que d'écrire une grammaire pour un décodeur d'octets. Il s'est avéré rapidement corriger l'implémentation de certaines instructions complexes de x86 (et pas seulement), grâce à un langage intermédiaire très pratique, un seul milieu (optimiseur), 2 backends (principalement C; en alternative - Java / C #, plus comme ce dernier, c'est-à-dire K. goto apparaît de temps en temps, mais pas étiqueté break).

Dans les articles suivants, il peut également y avoir des interfaces pour les langages gérés tels que DXBC, SPIR-V, ils utiliseront le backend Java / C #. Mais jusqu'à présent, seuls les codes machine sont dans les plans. les bytecodes nécessitent une approche spéciale. Ghidra Aide



Projet : Pcode Sleigh