Parallel Thread Execution (PTX oder NVPTX[1]) ist eine Befehlssatzarchitektur und virtuelle Maschine für die Ausführung von parallelen Threads, die in der CUDA-Programmierung von Nvidia verwendet wird. Der NVCC-Kompilierer übersetzt Code, der in CUDA, einer C++-ähnlichen Sprache, geschrieben wurde, in PTX-Anweisungen (eine Maschinensprache dargestellt als ASCII-Text). Der Grafiktreiber enthält einen Kompilierer, der die PTX-Anweisungen in den ausführbaren Binärcode übersetzt[2], der auf den Prozessorkernen der Nvidia-GPUs ausgeführt werden kann. Die GNU-Compiler-Kollektion verfügt auch über grundlegende Fähigkeiten zur PTX-Erzeugung im Zusammenhang mit OpenMP.[3] Inline-PTX-Assembly kann in CUDA verwendet werden.[4]
PTX verwendet einen beliebig großen Registersatz; die Ausgabe des Compilers erfolgt fast ausschließlich in Form von Einzelzuweisungen, wobei sich aufeinanderfolgende Zeilen im Allgemeinen auf aufeinanderfolgende Register beziehen. Programme beginnen mit Deklarationen der Form
.reg .u32 %r<335>; // deklariere 335 Register %r0, %r1, ..., %r334 vom Typ vorzeichenlosem 32-Bit Integer
Es handelt sich um eine Assemblersprache mit drei Argumenten. Fast alle Befehle geben ausdrücklich den Datentyp (in Form von Vorzeichen und Breite) an, mit dem sie arbeiten. Registernamen wird ein %-Zeichen vorangestellt und Konstanten sind Literale z. B.:
shr.u64 %rd14, %rd12, 32; // schiebe einen vorzeichenlosen 64-Bit Integer aus %rd12 nach rechts um 32 Positionen, Ergebnis in %rd14
cvt.u64.u32 %rd142, %r112; // konvertiere einen vorzeichenlosen 32-Bit Integer zu 64-Bit
Es gibt Prädikatsregister, aber kompilierter Code im Shader-Modell 1.0 verwendet diese nur in Verbindung mit Verzweigungsbefehlen; die bedingte Verzweigung ist
@%p14 bra $label; // gehe zu $label
Die Anweisung setp.cc.type setzt ein Prädikatsregister auf das Ergebnis des Vergleichs zweier Register des entsprechenden Typs. Es gibt auch eine Anweisung set, wobei set.le.u32.u64 %r101, %rd12, %rd28 das 32-Bit-Register %r101 zu 0xffffffff setzt, wenn das 64-Bit-Register %rd12 kleiner oder gleich dem 64-Bit-Register %rd28 ist. Andernfalls wird %r101 zu 0x00000000 gesetzt.
Es gibt ein paar vordefinierte Bezeichner, die Pseudoregister bezeichnen. Unter anderem enthalten %tid, %ntid, %ctaid, und %nctaid Thread-Indizes, Block-Dimensionen, Block-Indizes und Grid-Dimensionen.[5]
Die Befehle Laden (ld) und Speichern (st) beziehen sich auf einen von mehreren verschiedenen Zustandsbereichen, z. B. ld.param.
Es gibt acht Zustandsbereiche:[5]
.reg : Register.sreg : spezielle, plattformspezifische Register.const : geteilter, nur lesender Speicher.global : globaler Speicher von allen Threads.local : thread-lokaler Speicher.param : Parameter, die an der Kernel übergeben werden.shared : gemeinsamer Speicher von Threads im selben Block.tex : globaler Texturspeicher (veraltet)Geteilter Speicher wird in der PTX-Datei wie folgt deklariert:
.shared .align 8 .b8 pbatch_cache[15744]; // definiere 15.744 Bytes, ausgerichtet an 8-Byte-Grenzen
Das Schreiben von Kerneln in PTX erfordert die explizite Registrierung von PTX-Modulen über die CUDA-Treiber-API, was in der Regel umständlicher ist als die Verwendung der CUDA-Runtime-API und des CUDA-Compilers nvcc von Nvidia. Das GPU Ocelot-Projekt bot eine API zur Registrierung von PTX-Modulen neben CUDA-Runtime-API-Kernelaufrufen, allerdings wird GPU Ocelot nicht mehr aktiv weiterentwickelt.[6]