„Parallel Thread Execution“ – Versionsunterschied
| [ungesichtete Version] | [gesichtete Version] |
Inhalt gelöscht Inhalt hinzugefügt
CE |
Erstellt durch Übersetzen der Seite „Parallel Thread Execution“ |
||
| Zeile 1: | Zeile 1: | ||
| ⚫ | '''Parallel Thread Execution''' ('''PTX''' oder '''NVPTX'''<ref>{{Cite web|url=https://llvm.org/docs/NVPTXUsage.html|title=User Guide for NVPTX Back-end — LLVM 7 documentation|work=llvm.org}}</ref>) ist eine [[Befehlssatzarchitektur]] und [[virtuelle Maschine]] für die Ausführung von [[Parallelrechner|parallelen]] [[Thread (Informatik)|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 [[American Standard Code for Information Interchange|ASCII]]-Text). Der Grafiktreiber enthält einen Kompilierer, der die PTX-Anweisungen in den ausführbaren Binärcode übersetzt<ref>{{Cite web|url=https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html#instruction-set-ref|title=CUDA Binary Utilities|date=|accessdate=2019-10-19|last=|first=|work=docs.nvidia.com|language=en-us|archiveurl=|archivedate=|url-status=live}}</ref>, der auf den Prozessorkernen der [[Liste der Nvidia-Grafikprozessoren|Nvidia-GPUs]] ausgeführt werden kann. Die [[GNU Compiler Collection|GNU-Compiler-Kollektion]] verfügt auch über grundlegende Fähigkeiten zur PTX-Erzeugung im Zusammenhang mit [[OpenMP]].<ref>{{Cite web|url=https://gcc.gnu.org/wiki/nvptx|title=nvptx|work=GCC Wiki}}</ref> Inline-PTX-Assembly kann in CUDA verwendet werden.<ref>{{Cite web|url=http://docs.nvidia.com/cuda/inline-ptx-assembly/index.html|title=Inline PTX Assembly in CUDA|accessdate=2019-11-03|work=docs.nvidia.com|language=en-us}}</ref> |
||
{{Primary sources|date=August 2020}} |
|||
{{short description|Low-level parallel thread execution virtual machine and instruction set architecture}} |
|||
| ⚫ | '''Parallel Thread Execution''' ('''PTX''' |
||
== |
== Register == |
||
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<syntaxhighlight lang="asm"> |
|||
PTX uses an arbitrarily large register set; the output from the compiler is almost pure [[Static_single_assignment_form|single-assignment form]], with consecutive lines generally referring to consecutive registers. Programs start with declarations of the form |
|||
| ⚫ | |||
<syntaxhighlight lang="asm"> |
|||
</syntaxhighlight>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.:<syntaxhighlight lang="asm"> |
|||
| ⚫ | |||
| ⚫ | |||
</syntaxhighlight> |
|||
| ⚫ | |||
</syntaxhighlight>Es gibt Prädikatsregister, aber kompilierter Code im Shader-Modell 1.0 verwendet diese nur in Verbindung mit Verzweigungsbefehlen; die bedingte Verzweigung ist<syntaxhighlight lang="asm"> |
|||
| ⚫ | |||
| ⚫ | </syntaxhighlight>Die Anweisung <code>setp.cc.type</code> setzt ein Prädikatsregister auf das Ergebnis des Vergleichs zweier Register des entsprechenden Types. Es gibt auch einen Anweisung <code>set</code>, wobei <syntaxhighlight lang="asm" inline="">set.le.u32.u64 %r101, %rd12, %rd28</syntaxhighlight> das 32-Bit-Register <code>%r101</code> zu <code>0xffffffff</code> setzt, wenn das 64-Bit-Register <code>%rd12</code> kleiner oder gleich dem 64-Bit-Register <code>%rd28</code> ist. Andernfalls wird <code>%r101</code> zu <code>0x00000000</code> gesetzt. |
||
| ⚫ | Es gibt ein paar vordefinierte Bezeichner, die Pseudoregister bezeichnen. Unter anderem enthalten <code>%tid, %ntid, %ctaid</code>, und <code>%nctaid</code> Thread-Indizes, Block-Dimensionen, Block-Indizes und Grid-Dimensionen.<ref name="ptx-isa">{{Cite web|url=http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/ptx_isa_2.3.pdf|title=PTX ISA Version 2.3|publisher=}}</ref> |
||
It is a three-argument assembly language, and almost all instructions explicitly list the data type (in terms of sign and width) on which they operate. Register names are preceded with a % character and constants are literal, e.g.: |
|||
<syntaxhighlight lang="asm"> |
|||
| ⚫ | |||
| ⚫ | |||
</syntaxhighlight> |
|||
== Zustandsbereiche == |
|||
There are predicate registers, but compiled code in shader model 1.0 uses these only in conjunction with branch commands; the conditional branch is |
|||
| ⚫ | |||
<syntaxhighlight lang="asm"> |
|||
Es gibt acht Zustandsbereiche:<ref name="ptx-isa">{{Cite web|url=http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/ptx_isa_2.3.pdf|title=PTX ISA Version 2.3|publisher=}}<cite class="citation web cs1" data-ve-ignore="true">[http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/ptx_isa_2.3.pdf "PTX ISA Version 2.3"] <span class="cs1-format">(PDF)</span>.</cite></ref> |
|||
| ⚫ | |||
</syntaxhighlight> |
|||
| ⚫ | |||
| ⚫ | |||
* <code>.sreg</code> : spezielle, plattformspezifische Register |
|||
| ⚫ | |||
| ⚫ | |||
* <code>.local</code> : thread-lokaler Speicher |
|||
* <code>.param</code> : Parameter, die an der Kernel übergeben werden |
|||
| ⚫ | |||
| ⚫ | |||
Geteilter Speicher wird in der PTX-Datei wie folgt deklariert:<syntaxhighlight lang="nasm"> |
|||
| ⚫ | |||
| ⚫ | |||
</syntaxhighlight>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.<ref>{{Cite web|url=https://github.com/gtcasl/gpuocelot|title=GPUOCelot: A dynamic compilation framework for PTX|date=|accessdate=|last=|first=|work=github.com|archiveurl=|archivedate=|url-status=live}}</ref> |
|||
== |
== Siehe auch == |
||
| ⚫ | |||
There are eight state spaces:<ref name="ptx-isa"/> |
|||
| ⚫ | |||
* <code>.sreg</code> : special, read-only, platform-specific registers |
|||
| ⚫ | |||
| ⚫ | |||
* <code>.local</code> : local memory, private to each thread |
|||
* <code>.param</code> : parameters passed to the kernel |
|||
| ⚫ | |||
| ⚫ | |||
| ⚫ | |||
Shared memory is declared in the PTX file via lines at the start of the form: |
|||
| ⚫ | |||
<syntaxhighlight lang="nasm"> |
|||
| ⚫ | |||
</syntaxhighlight> |
|||
<!-- mov.u64 %rd9, pbatch_cache; |
|||
Shared memory is generally addressed via a kernel-global pointer set up at the start of the kernel by |
|||
--> |
|||
== Einzelnachweise == |
|||
Writing kernels in PTX requires explicitly registering PTX modules via the CUDA Driver API, typically more cumbersome than using the CUDA Runtime API and Nvidia's CUDA compiler, nvcc. The GPU Ocelot project provided an API to register PTX modules alongside CUDA Runtime API kernel invocations, though the GPU Ocelot is no longer actively maintained.<ref>{{cite web|url=https://github.com/gtcasl/gpuocelot|title=GPUOCelot: A dynamic compilation framework for PTX|last=|first=|date=|website=github.com|url-status=live|archive-url=|archive-date=|access-date=}}</ref> |
|||
== |
== Weblinks == |
||
| ⚫ | |||
| ⚫ | |||
| ⚫ | |||
== References == |
|||
| ⚫ | |||
{{Reflist}} |
|||
== External links == |
|||
| ⚫ | |||
| ⚫ | |||
Version vom 31. Juli 2022, 18:47 Uhr
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]
Register
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 vorzeichenlose 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 Types. Es gibt auch einen 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]
Zustandsbereiche
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]
Siehe auch
- Standard Portable Intermediate Representation (SPIR)
- CUDA binary (cubin) – a type of fat binary
Einzelnachweise
Weblinks
- ↑ User Guide for NVPTX Back-end — LLVM 7 documentation. In: llvm.org.
- ↑ CUDA Binary Utilities. In: docs.nvidia.com. Abgerufen am 19. Oktober 2019 (amerikanisches Englisch).
- ↑ nvptx. In: GCC Wiki.
- ↑ Inline PTX Assembly in CUDA. In: docs.nvidia.com. Abgerufen am 3. November 2019 (amerikanisches Englisch).
- ↑ a b PTX ISA Version 2.3. Referenzfehler: Ungültiges
<ref>-Tag. Der Name „ptx-isa“ wurde mehrere Male mit einem unterschiedlichen Inhalt definiert. - ↑ GPUOCelot: A dynamic compilation framework for PTX. In: github.com.