Koprozessoren
General Purpose Programmierung auf Grafikprozessoren
Stefan Zellmann
Lehrstuhl f¨ur Informatik, Universit¨at zu K¨oln
SS2019
Spracherweiterung / Aufrufsyntax Kernels
CUDA Compiler erweitert C++ um spezielle Syntax, um Compute Kernels aufzurufen.
_ _ g l o b a l _ _ v o i d k e r n e l () {
}
int m a i n () {
kernel < < < n u m _ b l o c k s , n u m _ t h r e a d s > > >();
}
Kernel ist Einstiegspunkt des GPU Programms (¨ahnlich wie main()f¨ur Host).
Programmausf¨uhrung kann von dort weiter verzweigen.
Globale Funktionkernel()wird von (#Bl¨ocke ×#Threads pro Block) Threads ausgef¨uhrt.
1D Grid, 1 Block, N Threads
_ _ g l o b a l _ _ v o i d s a x p y (f l o a t a , f l o a t* x , f l o a t* y ) {
int i = t h r e a d I d x . x ; y [ i ] = a * x [ i ] + y [ i ];
}
int m a i n () {
...
saxpy < < <1 , N > > >( a , x , y );
...
}
2D Grid, 1 Block, N× N Threads
_ _ g l o b a l _ _ v o i d m a t r i x _ a d d (f l o a t** a , f l o a t** b , f l o a t** c ) {
int i = t h r e a d I d x . x ; int j = t h r e a d I d x . y ;
c [ i ][ j ] = a [ i ][ j ] + b [ i ][ j ];
}
int m a i n () {
...
d i m 3 n u m _ t h r e a d s ( N , N );
m a t r i x _ a d d < < <1 , n u m _ t h r e a d s >( a , b , c );
...
}
2D Grid unterteilt in Bl¨ocke
_ _ g l o b a l _ _ v o i d m a t r i x _ a d d (f l o a t** a , f l o a t** b , f l o a t** c ) {
int i = b l o c k I d x . x * b l o c k D i m . x + t h r e a d I d x . x ; int j = b l o c k I d x . y * b l o c k D i m . y + t h r e a d I d x . y ; c [ i ][ j ] = a [ i ][ j ] + b [ i ][ j ];
}
int m a i n () {
...
// N /32 B l o e c k e in x - R i c h t u n g , M /32 in y - R i c h t u n g d i m 3 t h r e a d s _ p e r _ b l o c k (32 , 3 2 ) ;
d i m 3 b l o c k s ( r o u n d _ u p ( N / t h r e a d s _ p e r _ b l o c k . x ) , r o u n d _ u p ( M / t h r e a d s _ p e r _ b l o c k . y ));
m a t r i x _ a d d < < < blocks , t h r e a d s _ p e r _ b l o c k >( a , x , y );
...
}
Optimale Blockgr¨oße
Das ermitteln der optimalen Blockgr¨oße h¨angt von Registerzahl ab, die Kernel ben¨otigt.
Ubersetze Kernel mit¨ nvcc Option -Xptxas -v:
Kepler GK110, Maxwell, Pascal: 64k 32-bit Register File, max. 255 Register pro Thread.
Optimale Blockgr¨oße
Bestimme Blockgr¨oße mit CUDA Occupancy Calculator.xls
Threads identifizieren
In CUDA Kernels stehen eingebaute Variablen gridDim.{x|y|z}
blockIdx.{x|y|z}
blockDim.{x|y|z}
threadIdx.{x|y|z}
zur Verf¨ugung. gridDim gibt Anzahl der Bl¨ocke an. ¨Uber blockIdxwird Block identifiziert, ¨uber blockDimdessen Dimensionen und ¨uber threadIdxder lokaleThread Index im Block.
Der globale Thread Index kann mittels
u n s i g n e d i = b l o c k I d x . x * b l o c k D i m . x + t h r e a d I d x . x ; u n s i g n e d j = b l o c k I d x . y * b l o c k D i m . y + t h r e a d I d x . y ; u n s i g n e d k = b l o c k I d x . z * b l o c k D i m . z + t h r e a d I d x . z ;
ausgerechnet werden.
I Single Source Modell: Host Code und Kernel Code k¨onnen in der gleichen Compilation Unit aufgef¨uhrt werden, teilen sich
#includeDirektiven etc.
I vgl. “Kompilieren mit nvcc”, Code wird sp¨ater von CUDA Compiler aufgeteilt.
I Reduzierter C++ Sprachumfang in Kernels (Speicherallokation, Rekursion etc. nicht oder nur eingeschr¨ankt unterst¨utzt).
I Funktionen k¨onnen “wiederverwendet” werden: host &
device Funktionen.
Funktionsannotation
// F u n c t i o n can o n l y be u s e d on CPU f l o a t dot ( v e c 3 u , v e c 3 v ) {
r e t u r n u . x * v . x + u . y * v . y + u . z * v . z ; }
// F u n c t i o n can o n l y be u s e d on CPU _ _ h o s t _ _ f l o a t dot ( v e c 3 u , v e c 3 v ) {
r e t u r n u . x * v . x + u . y * v . y + u . z * v . z ; }
// F u n c t i o n can o n l y be u s e d on GPU _ _ d e v i c e _ _ f l o a t dot ( v e c 3 u , v e c 3 v ) {
r e t u r n u . x * v . x + u . y * v . y + u . z * v . z ; }
// F u n c t i o n can be u s e d on CPU and GPU
_ _ h o s t _ _ _ _ d e v i c e _ _ f l o a t dot ( v e c 3 u , v e c 3 v ) { r e t u r n u . x * v . x + u . y * v . y + u . z * v . z ; }
ausgerechnet werden.
Inlining
Herk¨ommlicher Funktionsaufruf (z. B. CPU) ⇒ Funktions-Stack, Argumente etc. werden in speziellen Registern gespeichert, werden, wenn Funktion zur¨uckkehrt, wieder freigegeben.
GPU und notorischer Registermangel⇒Funktionen fast immer inline.
Daher schlechter bis kein Support f¨ur Rekursion auf GPUs (w¨urde zu unkontrollierbarer Rekursionstiefe f¨uhren).
Generelle Empfehlungen
1.) Vermeide Divergenz, da Threads in Warp bei dynamischem Branching aufeinander warten.
2.) Koaleszierende Speicherzugriffe: Threads sollten immer von alignierten Speicheradressen lesen.
3.)Wegen Speicherzugriffslatenz: oft lohnt es sich, Berechungen immer wieder zuwiederholen,anstattErgebnisse
zwischenzuspeichern.
4.)Einmal auf der GPU, vermeide Kommunikation mit Host und f¨uhrem¨oglichst viele Berechnungen in Kernel aus.
Neben Kernels, die auf der GPU ausgef¨uhrt werden, steht Runtime Bibliothek zur Verf¨ugung, die das Interface zwischen Host und Device steuert.
CUDA Runtime Funktionen haben Pr¨afixcuda, z. B.
cudaMalloc(),cudaMemcpy()etc.
Runtime Initialisierung
Runtime wirdimplizit initialisiert. Erster CUDA Funktionsaufruf initialisiert Runtime.
I Runtime erstellt CUDA Kontextf¨ur jede installierte, CUDA-kompatible GPU.
I JIT Compilation und Laden von Device Code in GPU Speicher bei Kontexterzeugung.
I Kontext wird von allen CPU-Threads geteilt!
I cudaDeviceReset() zerst¨ort aktuellen Kontext, n¨achster Runtime Funktionsaufruf erstellt neuen Kontext.
Die meisten CUDA Runtime Funktionen geben Fehlercode zur¨uck:
c u d a E r r o r _ t err = c u d a G e t D e v i c e C o u n t ( . . . ) ; if ( err != c u d a S u c c e s s ) {
...
}
cudaGetLastError()undcudaPeekLastError(): pr¨ufe ob Kernel Fehler ausgel¨ost hat.
// Get l a s t error , r e s e t to c u d a S u c c e s s c u d a E r r o r _ t err = c u d a G e t L a s t E r r o r ();
// P e e k l a s t e r r o r w / o r e s e t
c u d a E r r o r _ t err = c u d a P e e k L a s t E r r o r ();
Achtung, Kernels werdenasynchron ausgef¨uhrt.LastError bezieht sich wom¨oglich nicht auf den richtigen Kernel⇒ f¨uge Synchronisation / Barrier ein:
// C a l l k e r n e l
kernel < < <1 , N > > >( p a r a m s );
// E r r o r s due to i n v a l i d k e r n e l c o n f i g u r a t i o n ( < < <... > > >) c u d a E r r o r _ t e r r 1 = c u d a G e t L a s t E r r o r ();
// K e r n e l e r r o r s ( e . g . out - of - b o u n d s m e m o r y a c c e s s ) c u d a E r r o r _ t err = c u d a D e v i c e S y n c h r o n i z e ();
Derartig synchronisierte Kernels laufen nicht mehr asynchron.
Fehlerbehandlung sollte nur im Debug Mode durchgef¨uhrt werden.
cudaDeviceSynchronize()ist generelles Synchronisationsprimitiv (falls Applikation Synchronisation erfordert).
cudaMalloc und cudaFree
I Wie in C/C++ muss DDR Speicher reserviert und sp¨ater wieder freigegeben werden.
I Device Zeiger: deklariere C-style raw pointer, weise mit cudaMalloc Adresse aus GPU Adressraum zu. Dieser Pointer steht auf dem Host zur Verf¨ugung, kann aber nur auf dem Device dereferenziert werden.
int* d _ p o i n t e r ;
c o n s t e x p r int N = 32;
c u d a M a l l o c (& d _ p o i n t e r , s i z e o f(int) * N );
...
c u d a F r e e ( d _ p o i n t e r );
I Bidirektionale Speichertransfers zwischen Host / Device (VRAM) und Speichertransfers zwischen VRAM
Speicherbereichen:cudaMemcpy().
I Aufruf Semantik wie ANSI-Cmemcpy, mit viertem Parameter, der Richtung angibt.
I enum cudaMemcpyKind:cudaMemcpyHostToHost(wie memcpy()),cudaMemcpyHostToDevice,
cudaMemcpyDeviceToHost,cudaMemcpyDeviceToDevice.
I VRAM wird mit Device Zeiger referenziert.
int h _ a r r a y = { 0 , 1 , 2 , 3 , 4 };
int* d _ a r r a y ;
c u d a M a l l o c (& d_array , s i z e o f(int) * 5);
c u d a M e m c p y ( d_array , h_array , s i z e o f(int) * 5 , c u d a M e m c p y H o s t T o D e v i c e );
CUDA Saxpy (1/2)
_ _ g l o b a l _ _ v o i d s a x p y (f l o a t a , f l o a t* x , f l o a t* y ) {
int i = t h r e a d I d x . x ; y [ i ] = a * x [ i ] + y [ i ];
}
int m a i n () {
c o n s t e x p r int N = ..;
f l o a t a ;
f l o a t * h_x , * h_y ;
// ... i n i t a , h_x , and h_y f l o a t * d_x , * d_y ;
s i z e _ t S = N * s i z e o f(f l o a t);
c u d a M a l l o c (& d_x , S );
c u d a M a l l o c (& d_y , S );
...
}
CUDA Saxpy (2/2)
...
c u d a M e m c p y ( d_x , h_x , S , c u d a M e m c p y H o s t T o D e v i c e );
c u d a M e m c p y ( d_y , h_y , S , c u d a M e m c p y H o s t T o D e v i c e );
saxpy < < <1 , N > > >( a , d_x , d_y );
// r e a d b a c k r e s u l t
c u d a M e m c p y ( h_y , d_y , S , c u d a M e m c p y D e v i c e T o H o s t );
c u d a F r e e ( d_x ); c u d a F r e e ( d_y );
...
}
I Steht limitiert (16 kb, 64 kb etc.) allen gemeinsam ausgef¨uhrten Threads auf SM zur Verf¨ugung.
I Keine direkte Adressierung via Zeiger, sondern spezielle Syntax in Kernel (CUDA Keyword shared ).
I Zwei Arten von Allokation: statischund dynamisch.
Statische Allokation
_ _ g l o b a l _ _ v o i d k e r n e l (int* da t a ) {
// S t a t i c s i z e s h a r e d m e m o r y _ _ s h a r e d _ _ int s h a r e d _ i n t s [ 6 4 ] ; // A c c e s s w i t h l o c a l t h r e a d ID s h a r e d _ i n t s [ t h r e a d I d x . x ]
= d a t a [ b l o c k I d x . x * b l o c k D i m . x + t h r e a d I d x . x ];
// A c c e s s to s h a r e d m e m o r y m u s t be s y n c h r o n i z e d _ _ s y n c t h r e a d s ();
// Now a c c e s s low - l a t e n c y m e m o r y ...
}
Dynamische Allokation
_ _ g l o b a l _ _ v o i d k e r n e l (int* da t a ) {
// S h a r e d memory , don ’ t s p e c i f y s i z e _ _ s h a r e d _ _ int s h a r e d _ i n t s [];
}
v o i d c a l l _ k e r n e l () {
// S p e c i f y v a r i a b l e s h a r e d m e m o r y s i z e // w h e n c a l l i n g k e r n e l ( m u s t s t i l l a d h e r e // to p l a t f o r m l i m i t s !)
kernel < < <
n u m _ b l o c k s ,
t h r e a d s _ p e r _ b l o c k s , s h a r e d _ m e m o r y _ s i z e // < < <
> > >( d a t a );
}
I Shared Memory lohnt sich, wenn Threads aus einer Warp h¨aufig auf Daten der anderen Threads in der Warp zugreifen.
I Je nach Architektur globaler Speicher gecached⇒ f¨ur trivialparallele Applikationen mit einfachen
Speicherzugriffsmustern lohnt sich Shared Memory u. Umst.
nicht.
I Matrix Operationen (z. B. DGEMM) profitieren eher von Shared Memory
I KopiereBlock in Shared Memory und bearbeite lokal.
Matrix Multiplikation mit Shared Memory C =A×B (vgl. NVIDIA Programming Guide.)
int row = t h r e a d I d x . y , col = t h r e a d I d x . x ; f l o a t cv = 0.0 f ;
for (int m = 0; m < ( A . w i d t h / B L O C K _ S I Z E ); ++ m ) { M a t r i x A s u b = G e t S u b M a t r i x ( A , b l o c k I d x . x , m );
M a t r i x B s u b = G e t S u b M a t r i x ( B , m , b l o c k I d x . y );
_ _ s h a r e d _ _ f l o a t As [ B L O C K _ S I Z E ][ B L O C K _ S I Z E ];
_ _ s h a r e d _ _ f l o a t Bs [ B L O C K _ S I Z E ][ B L O C K _ S I Z E ];
As [ row ][ col ] = A s u b [ row ][ col ];
Bs [ row ][ col ] = B s u b [ row ][ col ];
_ _ s y n c t h r e a d s ();
for (int e = 0; e < B L O C K _ S I Z E ; ++ e ) cv += As [ row ][ e ] * Bs [ e ][ col ];
_ _ s y n c t h r e a d s ();
}
C [ row ][ col ] = cv ;
Matrix Multiplikation mit Shared Memory
I 2D Grid, Idx.xf¨ur Zeile in Matrix A,Idx.y f¨ur Spalte in Matrix B.
I Iteriere ¨uber Blockgr¨oße, kopiere Submatrizen f¨ur ganzen Block in Shared Memory.
I Skalarproduktionoperation wird von jedem Thread in Shared Memory durchgef¨uhrt.
I Vor und nach Skalarprodukt m¨ussen Threads synchronisiert werden, damit nicht schon andere Threads aus Warp anderen Block in Shared Memory laden.
Matrix Multiplikation mit Shared Memory
Asub0 Asub1 Bsub1
Bsub0
BLOCK_SIZE BLOCK_SIZE
BLOCK_SIZE BLOCK_SIZE BLOCK_SIZE
BLOCK_SIZEBLOCK_SIZE
Global Thread ID y
Global Thread ID x
A B
C
Warp
Warp
I Read-Only Speicherbereich.
I Dedizierter Cache nur f¨ur Texturen.
I Dedizierte Hardware f¨ur Filtering (lineare Interpolation).
I Texturspeicher nicht direkt adressierbar, Verwaltung ¨uber spezielle Texturobjekte.
I Seit Nvidia Kepler Architektur (2012): Bindless Texturen und Texture Objects
I Texturen k¨onnen relativ flexibel erzeugt werden und m¨ussen nicht mehr explizit an Texture Units gebunden werden.
I Vorher: fixe Anzahl an Texturen im Programm.
I Texturspeicher ist nicht direkt adressierbar (keine Zeiger), stattdessen Zugriffsobjekt.
I IHVs halten Organisation des Texturspeichers geheim.
Texture Objects
Erzeuge Texture Object in Host Code:
c u d a T e x t u r e O b j e c t _ t obj ;
c u d a C r e a t e T e x t u r e O b j e c t (& obj , . . . ) ;
Dedizierte Texturzugriffsfunktionen in Kernel:
_ _ g l o b a l _ _ v o i d k e r n e l ( c u d a T e x t u r e O b j e c t _ t obj1D , c u d a T e x t u r e O b j e c t _ t obj2D , c u d a T e x t u r e O b j e c t _ t o b j 3 D ) {
f l o a t t1 = t e x 1 D ( obj1D , 0.5 f );
f l o a t t2 = t e x 2 D ( obj2D , m a k e _ f l o a t 2 ( 0 . 3 f , 0.4 f ));
f l o a t t3 = t e x 3 D ( obj3D ,
m a k e _ f l o a t 3 ( 0 . 3 f , 0.4 f , 0.5 f ));
}
⇒man kann Texturen nicht direkt adressieren.
Konfiguration von Texture Objects
Beim Erzeugen von Texture Objects legt man u. a. fest, I welche Dimensionalit¨at die Textur hat (1D, 2D, 3D), I welchen Datentyp die gespeicherten Texel haben, I wie interpoliert wird (n¨achster Nachbar; linear),
I ob Koordinaten normalisiert sind [0..1) oder nicht [0..Width), I sowie den Wrap Modus:
I Clamp I Wrap I Mirror I Border
