Wir schreiben eine Spezifikation für Nvidia Kepler (CUDA-Binärdateien, Sprachversion sm_30) für Ghidra

Für gängige Prozessorsprachen wurden bereits viele Spezifikationen für Ghidra geschrieben, für grafische jedoch nichts. Dies ist verständlich, da es seine eigene Spezifität hat: Prädikate, Konstanten, durch die Parameter übergeben werden, einschließlich anderer Dinge, die von Shadern geerbt wurden. Darüber hinaus ist das zum Speichern des Codes verwendete Format häufig proprietär, und Sie müssen es selbst umkehren.



In diesem Artikel werden wir anhand von zwei Beispielen herausfinden, was was ist.



Das erste Programm ist das einfachste Axpy (analog zu Hallo Welt für GPGPU). Die zweite hilft, die Implementierung von Bedingungen und Sprüngen auf der GPU zu verstehen, da dort ist alles anders.



Alle Nvidia-Sprachen verwenden die Little-Endian-Codierung. Kopieren Sie daher die Bytes aus dem Hex-Editor sofort in umgekehrter Reihenfolge von 8 Teilen in ein Notizbuch (z. B. Notepad ++) (die Länge der Anweisungen ist hier konstant). Dann übersetzen wir ihn über den Taschenrechner des Programmierers (der von Microsoft ist geeignet) in Binärcode. Als nächstes suchen wir nach Übereinstimmungen, setzen die Anweisungsmaske und dann die Operanden zusammen. Um die Maske zu dekodieren und zu suchen, wurden der Hex-Editor und der cuobjdump-Disassembler verwendet. Manchmal ist ein Assembler erforderlich, wie in AMDGPU (da der Disassembler dort nicht verfügbar ist, dies ist jedoch ein Thema für einen separaten Artikel). Das funktioniert so: Wir versuchen, alle verdächtigen Bits im Taschenrechner nacheinander zu invertieren, erhalten dann einen neuen Hexadezimalwert für die Bytes und ersetzen sie durch eine über nvcc kompilierte Binärdatei oder einen Assembler, falls vorhanden, was nicht immer der Fall ist.Dann checken wir durch cuobjdump.



Ich verteile die Quelle im Format (hauptsächlich in C, ohne Pluspunkte und OOP für eine engere Verbindung mit dem GPU-Code des Computers) und disasm + Bytes auf einmal, da es bequemer ist, müssen sie einfach nicht ausgetauscht werden.



Kopieren Sie es in axpy.cu und kompilieren Sie es über cmd: nvcc axpy.cu --cubin --gpu-architektur sm_30 Zerlegen Sie die

resultierende ELF-Datei mit dem Namen axpy.cubin an derselben Stelle: cuobjdump axpy.cubin -sass



Beispiel 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];
}

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

Es ist nicht schwer zu erraten, dass die Tests ursprünglich für Maschinencode entwickelt wurden, sodass der Compiler nichts zu optimieren hatte. Für alles andere müssten Sie die Optimierungen manuell abbrechen. In komplexen Beispielen ist dies möglicherweise überhaupt nicht möglich. In solchen Fällen müssen Sie sich daher auf den Dekompiler und das Frontend verlassen.



Im Allgemeinen gilt die Regel, dass wir zum Testen des Frontends jedes einfache (mit einem Minimum an möglichen Optimierungen) zuerst geeignete (Wiedergabefehler) Beispiel verwenden. Im Übrigen wird der dekompilierte Code bereits optimiert (oder durch Refactoring einfach irgendwie korrigiert). Aber im Moment ist die Hauptaufgabe zumindest genau der richtige Code, der dasselbe tut wie Maschinencode. Dies ist "Softwaremodellierung". "Softwaremodellierung" selbst bedeutet nicht Refactoring, Übersetzung von C nach C ++, Wiederherstellung von Klassen und vor allem Dinge wie die Identifizierung von Vorlagen.



Jetzt suchen wir nach Mustern für Mnemonik, Operanden und Modifikatoren.



Dazu vergleichen wir die Bits (in binärer Darstellung) zwischen verdächtigen Anweisungen (oder Zeichenfolgen, wenn es bequemer ist, sie so aufzurufen). Sie können auch das verwenden, was andere Benutzer in ihren Fragen zum Stackoverflow posten, z. B. "Hilf mir, Binär- / Sass- / Maschinencode zu verstehen", Tutorials (auch auf Chinesisch) und andere Ressourcen verwenden. Die Hauptoperationsnummer ist also in den Bits 58-63 gespeichert, aber es gibt auch zusätzliche Bits 0-4 (sie unterscheiden die Anweisungen "I2F", "ISETP", "MOV32I"), irgendwo anstelle von ihnen 0-2 (zur Vernachlässigung 3- 4 Bits in leeren Anweisungen, in der Spezifikation sind sie als "UNK" gekennzeichnet).



Bei Registern und konstanten Zahlen können Sie mit dem Disassembler experimentieren, um alle Bits zu finden, die sich auf die Dump-Ausgabe auswirken, wie das unter dem Spoiler angegebene. Alle Felder, die ich gefunden habe, befinden sich in der Spezifikation auf Github, Datei CUDA.slaspec, Abschnittstoken.



Dann müssen Sie Adressen für die Register finden, wieder sind sie auf Github. Dies ist notwendig, weil Auf der Mikroebene betrachtet Sleigh Register als globale Variablen in einem Raum vom Typ "register_space", aber seitdem Da ihr Leerzeichen nicht als "ableitbar" markiert ist (und dies sicherlich nicht sein kann), werden sie im Dekompiler entweder zu lokalen Variablen (meistens mit der "Var" -Schnittstelle, aber manchmal war auch das "lokale" Präfix ähnlich) oder zu Parametern (" param_ "). SP war nie nützlich, es wird hauptsächlich formal benötigt, um sicherzustellen, dass der Dekompiler funktioniert. Für die Emulation wird ein PC (etwa IP von x86) benötigt.



Dann gibt es Prädikatenregister, so etwas wie Flags, aber eher "allgemeiner Zweck" als für einen vorbestimmten Zweck, wie Überlauf, (Un-) Gleichheit mit Null usw.

Dann ein Sperrregister, um eine Reihe von Anweisungen ISCADD .CC und IMAD.HI zu simulieren, da Der erste von ihnen in meiner Implementierung berechnet für sich selbst und für den zweiten, um zu vermeiden, dass ein Teil der Summe auf die hohen 4 Bytes übertragen wird, weil Dies wird die Dekompilierung ruinieren. Dann müssen Sie jedoch das nächste Register sperren, bis der IMAD.HI-Vorgang abgeschlossen ist. Ähnliches, d.h. Die Diskrepanz zwischen der offiziellen Dokumentation und der erwarteten Ausgabe des Dekompilierers war bereits in der SPU für dieselbe Ghidra vorhanden.



Dann gibt es spezielle Register, die bisher über cpool implementiert wurden. In Zukunft plane ich, sie durch die standardmäßig definierten Zeichen für eine Art "ableitbaren" Raum zu ersetzen. Dies sind die gleichen threadIdx, blockIdx.



Dann binden wir die Variablen an die Felder dest, par0, par1, par2, res. Dann gibt es Untertabellen und danach - worum es ging - die Haupttabellen (Root-Tabellen) mit den Hauptanweisungen.



Hier müssen Sie sich strikt an das Format der "Mnemonik-Operanden" halten. Für Modifikatoren gilt jedoch eine Ausnahme, die jedoch an die Mnemonik oder an den Abschnitt mit Operanden angehängt werden muss. Es sind keine anderen Formate zulässig, auch derselbe Hexagon-DSP-Asm muss an diese Syntax angepasst werden, was jedoch nicht sehr schwierig ist.



In der letzten Phase wird die Implementierung für die Anweisungen in der Pcode-Mikroprogrammierungssprache geschrieben. Das einzige, was ich aus dem ersten Beispiel heraus erwähnen möchte, sind die Anweisungen ISCADD .CC und IMAD.HI, bei denen der erste einen Zeiger auf Register nimmt und diese als Zeiger für 8 Bytes anstelle von 4 dereferenziert. Dies geschieht absichtlich, um sich besser an den Dekompiler anzupassen und sein Verhalten, trotz dessen, was in der Nvidia-Dokumentation über die Überweisung eines Teils des Betrags geschrieben steht.



Für das zweite Beispiel ist es besser, das Kontrollkästchen in den Dekompilereinstellungen neben der Inschrift "Prädikation vereinfachen" zu aktivieren. Der Punkt ist, dass Prädikate ein und dieselbe Bedingung für verschiedene Anweisungen sind, in der Tat nichts weiter als die bekannte "SIMD" oder ihr nächstes Analogon. Jene. Wenn das Prädikatbit gesetzt ist, werden die Anweisungen außerdem in einer Reihe ausgeführt.



Sie müssen sich auch angewöhnen, sofort eine Implementierung für jede Anweisung zu schreiben und nicht nur einen Prototyp (mnemonische Operanden), da es auch einen Dekompiler, einen Emulator und andere Analysatoren gibt.

Im Allgemeinen ist das Schreiben einer Implementierung in Pcode jedoch eine noch einfachere Aufgabe als das Schreiben einer Grammatik für einen Byte-Decoder. Es stellte sich schnell heraus, dass die Implementierung für einige komplexe Anweisungen von x86 (und nicht nur) dank einer sehr praktischen Zwischensprache, einem einzelnen Mittelland (Optimierer), 2 Backends (meistens C; als Alternative - Java / C #, eher wie letztere, d. H. (da von Zeit zu Zeit ein goto erscheint, aber keine beschriftete Pause).

In den folgenden Artikeln gibt es möglicherweise auch Frontends für verwaltete Sprachen wie DXBC, SPIR-V, die das Java / C # -Backend verwenden. Bisher sind jedoch nur Maschinencodes geplant. Bytecodes erfordern einen speziellen Ansatz. Ghidra Hilfe



Projekt : Pcode Sleigh