Execuție fir paralel - Parallel Thread Execution

Paralel Executie Subiect ( PTX sau NVPTX ) este un nivel scăzut- paralel fir de execuție mașină virtuală și arhitectura set de instrucțiuni utilizate în Nvidia e CUDA mediu de programare. NVCC compilator traduce codul scris în CUDA, un C ++ limbaj -cum ar fi, în instrucțiuni PTX, iar driverul grafic conține un compilator care traduce instrucțiunile PTX într - un cod binar , care poate fi rulat pe nuclee de procesare ale Nvidia GPU . Colecția GNU Compiler are, de asemenea, capacitatea de bază pentru generarea PTX în contextul descărcării OpenMP . Ansamblul PTX în linie poate fi utilizat în CUDA.

Registrele

PTX utilizează un set de registre arbitrar de mare; ieșirea din compilator este o formă aproape pură cu o singură atribuire , cu linii consecutive care se referă în general la registre consecutive. Programele încep cu declarațiile formularului 

.reg .u32 %r<335>;            // declare 335 registers %r0, %r1, ..., %r334 of type unsigned 32-bit integer

Este un limbaj de asamblare cu trei argumente și aproape toate instrucțiunile enumeră în mod explicit tipul de date (în ceea ce privește semnul și lățimea) pe care operează. Numele de registre sunt precedate de un caracter% și constantele sunt literal, de exemplu: 

shr.u64 %rd14, %rd12, 32;     // shift right an unsigned 64-bit integer from %rd12 by 32 positions, result in %rd14
cvt.u64.u32 %rd142, %r112;    // convert an unsigned 32-bit integer to 64-bit

Există registre de predicate, dar codul compilat în modelul shader 1.0 le utilizează numai împreună cu comenzi de ramură; ramura condițională este 

@%p14 bra $label;             // branch to $label

setp.cc.typeInstrucțiunea stabilește un registru predicat la rezultatul compararea a două registre de tip corespunzătoare, există , de asemenea , o setinstrucțiune, în cazul în care seturi registrul pe 32 de biți pentru a în cazul în care registrul pe 64 de biți este mai mică sau egală cu registrul pe 64 de biți . Altfel este setat la . set.le.u32.u64 %r101, %rd12, %rd28%r1010xffffffff%rd12%rd28%r1010x00000000

Există câțiva identificatori predefiniți care denotă pseudoregistori. Printre altele, %tid, %ntid, %ctaidși %nctaidconțin, respectiv, indici de fire, dimensiuni bloc, indici bloc și dimensiuni grilă.

Spații de stat

Comenzile Load ( ld) și store ( st) se referă la unul dintre mai multe spații de stare distincte (bănci de memorie), de ex ld.param. Există opt spații de stare:

  • .reg : registre
  • .sreg : registre speciale, numai în citire, specifice platformei
  • .const : memorie partajată, numai în citire
  • .global : memorie globală, partajată de toate firele
  • .local : memorie locală, privată pentru fiecare fir
  • .param : parametrii trecuți nucleului
  • .shared : memoria partajată între firele dintr-un bloc
  • .tex : memorie textură globală (depreciată)

Memoria partajată este declarată în fișierul PTX prin linii la începutul formularului: 

.shared .align 8 .b8 pbatch_cache[15744]; // define 15,744 bytes, aligned to an 8-byte boundary

Scrierea nucleelor ​​în PTX necesită înregistrarea în mod explicit a modulelor PTX prin intermediul API-ului CUDA Driver, de obicei mai greoaie decât utilizarea API-ului CUDA Runtime și a compilatorului CUDA Nvidia, nvcc. Proiectul GPU Ocelot a furnizat un API pentru a înregistra module PTX alături de invocațiile nucleului CUDA Runtime API, deși GPU Ocelot nu mai este întreținut în mod activ.

Vezi si

Referințe

linkuri externe