Parallel Thread Execution - Parallel Thread Execution

Parallel Thread Execution (PTXnebo NVPTX[1]) je nízká úroveň paralelní vlákno provedení virtuální stroj a architektura sady instrukcí použito v Nvidia je CUDA programovací prostředí. The nvcc překladač překládá kód napsaný v CUDA, a C ++ -jako jazyk, do pokynů PTX a grafický ovladač obsahuje kompilátor, který převádí pokyny PTX do binárního kódu[2] které lze spustit na procesorových jádrech GPU Nvidia. The Sbírka překladačů GNU má také základní schopnosti pro generování PTX v kontextu OpenMP vykládka.[3] V CUDA lze použít vloženou sestavu PTX.[4]

Registry

PTX používá libovolně velkou sadu registrů; výstup z kompilátoru je téměř čistý formulář pro jedno přiřazení, přičemž po sobě jdoucí řádky obecně odkazují na po sobě jdoucí registry. Programy začínají prohlášeními formuláře

.reg .u32 % r<335>; // deklarovat 335 registrů% r0,% r1, ...,% r334 typu 32bitové celé číslo bez znaménka

Jedná se o tříparametrový montážní jazyk a téměř všechny pokyny výslovně uvádějí datový typ (z hlediska znaménka a šířky), na kterém fungují. Před názvy registrů je uveden znak% a konstanty jsou doslovné, např .:

shr.u64 % rd14, % rd12, 32; // posunout vpravo nepodepsané 64bitové celé číslo z% rd12 o 32 pozic, výsledkem bude% rd14cvt.u64.u32 % rd142, % r112; // převést nepodepsané 32bitové celé číslo na 64bitové

Existují predikátové registry, ale kompilovaný kód v shaderovém modelu 1.0 je používá pouze ve spojení s příkazy větví; podmíněná větev je

@%p14 podprsenka $ štítek; // větev na $ label

The typ setp.cc. instrukce nastaví predikátový registr na výsledek porovnání dvou registrů příslušného typu, existuje také a soubor instrukce, kde set.le.u32.u64 % r101, % rd12, % rd28 nastavuje 32bitový registr % r101 na 0xffffffff pokud je 64bitový registr % rd12 je menší nebo roven 64bitovému registru % rd28. v opačném případě % r101 je nastaven na 0x00000000.

Existuje několik předdefinovaných identifikátorů, které označují pseudoregistery. Mezi ostatními, % tid,% ntid,% ctaid, a % nctaid obsahují indexy závitů, kóty bloků, indexy bloků a kóty mřížky.[5]

Stavové prostory

Načíst (ld) a uložte (Svatý) příkazy odkazují na jeden z několika odlišných stavových prostorů (paměťových bank), např. ld. paramExistuje osm stavových prostorů:[5]

  • .reg : registry
  • .reg : speciální registry určené pouze ke čtení a specifické pro platformu
  • .const : sdílená paměť pouze pro čtení
  • .globální : globální paměť sdílená všemi vlákny
  • .místní : místní paměť, soukromá pro každé vlákno
  • .param : parametry předané jádru
  • . sdílené : paměť sdílená mezi vlákny v bloku
  • .tex : globální paměť textury (zastaralé)

Sdílená paměť je v souboru PTX deklarována pomocí řádků na začátku formuláře:

. sdílené . zarovnat 8 .b8 pbatch_cache[15744]; // definuje 15 744 bajtů, zarovnáno na hranici 8 bajtů

Psaní jader v PTX vyžaduje výslovnou registraci modulů PTX prostřednictvím rozhraní CUDA Driver API, což je obvykle těžkopádnější než používání rozhraní CUDA Runtime API a kompilátoru NVIDIA CUDA, nvcc. Projekt GPU Ocelot poskytl API pro registraci modulů PTX vedle vyvolání jádra CUDA Runtime API, ačkoli GPU Ocelot již není aktivně udržován.[6]

Viz také

Reference

  1. ^ „Uživatelská příručka pro back-end NVPTX - dokumentace LLVM 7“. llvm.org.
  2. ^ „CUDA Binary Utilities“. docs.nvidia.com. Citováno 2019-10-19.
  3. ^ "nvptx". GCC Wiki.
  4. ^ „Inline PTX Assembly in CUDA“. docs.nvidia.com. Citováno 2019-11-03.
  5. ^ A b „PTX ISA verze 2.3“ (PDF).
  6. ^ „GPUOCelot: Dynamický kompilační rámec pro PTX“. github.com.

externí odkazy