Parallell trådutförande
Parallell Thread Execution ( PTX eller NVPTX ) är en lågnivå parallell trådexekvering virtuell maskin och instruktionsuppsättningsarkitektur som används i Nvidias CUDA - programmeringsmiljö . NVCC - kompilatorn översätter kod skriven i CUDA, ett C++ -liknande språk, till PTX-instruktioner (ett assemblerspråk representerat som ASCII -text), och grafikdrivrutinen innehåller en kompilator som översätter PTX-instruktionerna till den körbara binära koden som kan köras på bearbetningskärnorna i Nvidia GPU:er . GNU Compiler Collection har också grundläggande förmåga för PTX-generering i samband med OpenMP- avlastning. Inline PTX-montage kan användas i CUDA.
Register
PTX använder en godtyckligt stor registeruppsättning; utgången från kompilatorn är nästan ren enkeltilldelningsform , med på varandra följande rader som i allmänhet hänvisar till konsekutiva register. Program börjar med deklarationer av formuläret
.reg .u32 %r < 335 > ; // deklarera 335 register %r0, %r1, ..., %r334 av typen osignerat 32-bitars heltal
Det är ett sammansättningsspråk med tre argument, och nästan alla instruktioner listar uttryckligen den datatyp (i termer av tecken och bredd) som de verkar på. Registernamn föregås av ett %-tecken och konstanter är bokstavliga, t.ex.
shr.u64 %rd14 , %rd12 , 32 ; // flytta höger ett osignerat 64-bitars heltal från %rd12 med 32 positioner, resulterar i %rd14 cvt.u64.u32 %rd142 , %r112 ; // konvertera ett osignerat 32-bitars heltal till 64-bitars
Det finns predikatregister, men kompilerad kod i shader modell 1.0 använder dessa endast i samband med grenkommandon; den villkorliga grenen är
@% p14 bh $label ; // förgrena sig till $label
Instruktionen setp.cc.type sätter ett predikatregister till resultatet av att jämföra två register av lämplig typ, det finns också en uppsättningsinstruktion , där set.le.u32.u64 %r101 , %rd12 , %rd28 ställer in 32-bitars registrera %r101 till 0xffffffff om 64-bitarsregistret %rd12 är mindre än eller lika med 64-bitarsregistret %rd28 . Annars %r101 satt till 0x00000000 .
Det finns några fördefinierade identifierare som betecknar pseudoregister. Bland annat %tid, %ntid, %ctaid och %nctaid trådindex, blockdimensioner, blockindex och griddimensioner.
Statliga utrymmen
Ladda ( ld ) och lagra ( st ) kommandon hänvisar till ett av flera distinkta tillståndsutrymmen (minnesbanker), t.ex. ld.param . Det finns åtta delstatsutrymmen:
-
.reg: registrerar -
.sreg: särskilda, skrivskyddade, plattformsspecifika register -
.const: delat, skrivskyddat minne -
.global: globalt minne, delat av alla trådar -
.local: lokalt minne, privat för varje tråd -
.param: parametrar som skickas till kärnan -
.shared: minne som delas mellan trådar i ett block -
.tex: globalt texturminne (utfasad)
Delat minne deklareras i PTX-filen via rader i början av formuläret:
.shared .align 8 .b8 pbatch_cache [ 15744 ] ; // definiera 15 744 byte, justerade till en 8-byte gräns
Att skriva kärnor i PTX kräver explicit registrering av PTX-moduler via CUDA Driver API, vanligtvis mer besvärligt än att använda CUDA Runtime API och Nvidias CUDA-kompilator, nvcc. GPU Ocelot-projektet tillhandahöll ett API för att registrera PTX-moduler tillsammans med CUDA Runtime API-kärnanrop, även om GPU Ocelot inte längre aktivt underhålls.
Se även
- Standard Portable Intermediate Representation (SPIR)
- CUDA binär (kubin) – en typ av fettbinär