GPU-laitteisto ja tehokkaan ohjelmakoodin kirjoittaminen

Nvidian toteutus

Warpit

• Nvidian näytönohjaimet suorittavat ytimen ohjelmakoodia ns. warpeissa
• Yksi warp koostuu 32 säikeestä, joiden indeksinumerot ovat peräkkäisiä:

• Ohjelmoija jakaa siis säikeet säieryhmiin, jonka jälkeen GPU-laitteisto jakaa säieryhmät edelleen warpeihin
• Laskentayksikkö (streaming multiprocessor) sisältää useamman prosessointielementin (CUDA core) ja yhden tai useamman warppivuorottajan, jotka jakavat laskentayksikön resurssit warppien kesken:

• Jokaisella säikeellä on oma ohjelmalaskuri, rekisterit ja yksityinen muisti
• Jokainen säie suorittaa siis ytimen lähdekoodia itsenäisesti loogisella tasolla
• GPU-raudan mielessä samaan warppiin kuuluvat säikeet suoritetaan kuitenkin yhdessä
• Osa säikeitä kytketään automaattiseti pois käytöstä silloin kun warpiin kuuluvien säikeiden suorituspolut haarautuvat

Warppien vuorotus

• Laskentayksikölle allokoidut warpit muodostavat "warppipoolin"
• Warppipoolissa olevat warpit voivat olla kahdessa eri tilassa:
  • Valmiita (ready) eli ainakin yksi warppiin kuuluva säie on valmis suorittamaan komennon
  • Odottamassa (waiting) eli warpissa ei ole yhtään säiettä, joka olisi valmis suorittamaan seuraavan komennon
• Warppiin kuuluvat säikeet saattavat olla odottamassa esimerkiksi edeltävän komennon valmistumista
• Warppivuorottimet valitsevat warppipoolista joka kellojaksolla joukon warppeja (olettaen, että ainakin yksi on olemassa) ja asettavat valitut warpit suoritukseen prosessointielementeille (issued)

• Warppivuorotin voi valita suoritettavan käskyn:
  • Useamman valmiin warpin joukosta
  • Useamman toisistaan riippumattoman käskyn joukosta

Compute capability (CC) -numero

• Nvidia ilmoittaa valmistamiensa näytönohjaimien kyvykkyydet compute capability -numeron muodossa
• Tällä kurssilla käytössä olevien GPU:n CC-numerot ovat
  • Piraija (Nvidia Tesla K40c / Kepler): 3.5
  • Mako (Nvidia GeForce GTX580 / Fermi): 2.0
• CC-numero kertoo esimerkiksi GPU:n tukemat CUDA-ominaisuudet (ja vastaavat OpenCL-ominaisuudet) ja rautatason toteutuksen pääpiirteet

Compute capability 2.x

Laskentayksiköt

• CC 2.0 -laitteen laskentayksikkö sisältää:
  • 32 prosessointielementtiä (CUDA-ydintä)
  • 4 erikoisyksikköä yksinkertaisen tarkkuuden erikoisfunktioiden käsittelyyn
  • 8 PE:tä yhdistetään silloin kun laskentayksikkö suorittaa kaksinkertaisen tarkkuuden liukulukuoperaatioita
  • 64KB yhdistetty lokaali (jaettua) muisti ja L1-välimuisti
• CC 2.1 -laitteen laskentayksikkö sisältää:
  • 48 prosessointielementtiä
  • 8 erikoisyksikköä yksinkertaisen tarkkuuden erikoisfunktioiden käsittelyyn
  • 64KB yhdistetty lokaali (jaettua) muisti ja L1-välimuisti

• Sekä CC 2.0 ja CC 2.1 laitteet sisältävät 2 warppivuorottajaa
• Yhden kellojakson aikana säielohkovuorotin voi vuorottaa
  • yhden käskyn CC 2.0 -laitteilla ja
  • kaksi toisistaan riippumatonta käskyä CC 2.1 -laitteilla
• Yksi vuorotin voi vuorottaa warppeja vain puolelle prosessointielementeistä
• Ensimmäinen warppivuorotin vuorottaa kaikki parittomat warpit ja toinen kaikki parilliset warpit
• Tuplatarkkuuslaskennassa vain toinen warppivuorotin on aktiivinen

Globaali muisti

• Jokaisessa laskentayksikössä on oma L1-tason välimuisti ja kaikki laskentayksiköt jakavat yhteisen L2-tason välimuistin
• Samaan warppiin kuuluvat säikeet tekevät muistipyynnöt yhdessä!
• L1-välimuistin rivin leveys on 128 tavua (L1 voidaan kytkeä osittain pois päältä)
• Jos warppiin kuuluvien säikeiden muistipyynnöt eivät viittaa samalle 128 tavun välimuistiriville, jaetaan warpin muistipyyntö useaan 128-tavun osaan
• L2-tason välimuisti palvelee muistipyyntöjä 32 tavun paloissa

Lokaali muisti

• Lokaali muisti on jaettu 32 pankkiin siten, että peräkkäiset 32 bitin sanat kuuluvat peräkkäisin pankkeihin
• Kunkin pankki voi palvella 32-bittiä kahden kellojakson aikana
• Mikäli samaan warppiin kuuluvan säikeet viittaavat samaan pankkiin samanaikaisesti, syntyy ns. pankkikonflikti
• Pankkikonfliktit käsitellään siten, että samaan pankkiin osoittaneet pyynnöt käsitellään peräkkäin
• Samaan 32-bitin sanaan viittaavat pyynnöt eivät aiheuta pankkikonfliktia

Compute capability 3.x

Laskentayksiköt

• CC 3.x -laitteen laskentayksikkö sisältää:
  • 192 yksinkertaisen tarkkuuden PE:tä
  • 64 kaksinkertaisen tarkkuuden PE:tä (K40c)
  • 4 warppivuorottajaa
  • 64KB yhdistetty lokaali (jaettua) muisti ja L1-välimuisti
• Kukin warppivuorotin kykenee vuorottamaan kaksi toisistaan riippumatonta komentoa yhdestä valmiina olevasta warpista

Globaali muisti

• Globaaliin muistiin viittaavat muistipyynnöt kulkevat oletuksena vain L2-tason välimuistin läpi
• L2-tason välimuisti käyttäytyy pääpiirteittäin samoin kuin CC 2.x laitteiden L2-tason välimuisti

Lokaali muisti

• Lokaali muisti on jaettu pankkeihin kuten CC 2.x -laitteidenkin tapauksessa, mutta pankin "leveys" koko voidaan asettaa cudaDeviceSetSharedMemConfig()-aliohjelmalla joko 32 tai 64 bittiin
• Edellinen tarkoitaa siis sitä, että
  • Lokaali muisti on jaettu 32 pankkiin siten, että peräkkäiset 32 bitin sanat kuuluvat peräkkäisin pankkeihin tai
  • Lokaali muisti on jaettu 32 pankkiin siten, että peräkkäiset 64 bitin sanat kuuluvat peräkkäisin pankkeihin

Compute capability 5.x

Laskentayksiköt

• CC 5.x -laitteen laskentayksikkö sisältää:
  • 128 yksinkertaisen tarkkuuden PE:tä
  • Neljä kaksinkertaisen tarkkuuden PE:tä
  • 4 warppivuorottajaa
  • 24KB L1-välimuisti
  • 64KB lokaali (jaettu) muisti
• Kukin warppivuorotin kykenee vuorottamaan yhden komennon yhdestä valmiina olevasta warpista

Globaali muisti

• Globaaliin muistiin viittaavat muistipyynnöt kulkevat oletuksena vain L2-tason välimuistin läpi
• L2-tason välimuisti käyttäytyy pääpiirteittäin samoin kuin CC 2.x laitteiden L2-tason välimuisti

Lokaali muisti

• Lokaali muisti on jaettu 32 pankkiin siten, että peräkkäiset 32 bitin muistiosoitteen kuuluvat peräkkäisin pankkeihin
• Lokaali muisti käyttäytyy pääpiirteittäin samoin kuin CC 2.x -laitteissa

AMD:n toteutus

Wavefront

• AMD:n näytönohjaimet suorittavat ytimen ohjelmakoodia ns. wavefronteissa
• Yksi wavefront koostuu 64 säikeestä, joiden indeksinumerot ovat peräkkäisiä:

Laskentayksiköt

• GCN-näytönohjaimen laskentayksikkö sisältää:
  • Yhden skalaariprosessointiyksikön
  • Neljä 16 prosessointielementin mittaista vektoriyksikköä
  • L1-välimuistin
  • Lokaalin muistin / Local Data Share (LDS) (tyypillisesti 64KB)

Wavefronttien vuorotus

• Kaikki vektoriyksikön prosessointielementit suorittavat saman käskyn kellojaksolla
• Tyypillisesti vektoriyksikkö suorittaa samaan wavefrontiin liittyvän komennon neljän kellojakson aikana eli kaikki 64 säiettä suoritetaan effektiivisesti yhdessä
• Eri laskentayksiköt ja eri vektoriyksiköt laskentayksikön sisällä voivat suorittaa eri komennon samalla kellojaksolla

• Komentovirta voi sisältää skalaarikäskyjä skalaariyksikölle ja vektorikäskyjä vektoriyksiköille
• Yhden kellojakson aikana voidaan vuorottaa skalaarikäsky, yksi vektorikäsky, muistioperaatio ja haarautuvuuskäsky (branch operation)
• Yksi vektoriyksikkö suorittaa yhtä vektorikäskyä neljä kellojaksoa, joten vektorikäskyt vuorotetaan vuorotellen eri vektoriyksiköille

Prosessointielementtien käyttöasteesta

• Eri komentojen suorittaminen vaati eri määrän kellojaksoja:

• Käytännössä CPU:t ja GPU:t suorittavat käskyjä liukuhihnalla, jolloin joka kellojaksolla voidaan aloittaa uuden käskyn suorittaminen:

• Usein liukuhihnalle saattaa kuitenkin jäädä tyhjiä välejä eli ns. kuplia

• GPU:n tapauksessa tilanen on vielä pahempi, koska kupla koskettaa tällöin kokonaista warppia/wavefronttia:

• On siis tärkeää, että warppi/wavefrontvuorottaja voi asettaa komennon prosessointielementtien liukuhihnoille joka kellojaksolla
• Käytännössä tämä onnistuu siten, että säikeitä on huomattavasti enemmän kuin prosessointielementtejä
• Myöskin käskytason rinnakkaisuus auttaa asian suhteen

Warppien ja Wavefronttien vaikutus

• Warpit ja wavefrontit suorittavat ohjelmakoodia yhdessä, joten kaikki niihin kuuluvat säikeet käyvät läpi kaikki tarvittavat suorituspolut
• Alla kuvaaja, jossa x-akselilla on todennäköisyys, jolla yksi säie haarautuu ja y-akselilla todennäköisyys, että koko warp/wavefront joutuu haarautumaan:

• Ehtolauseilla voi siis olla erittäin merkittävä vaikutus ytimen suorituskykyyn
• Pyri siis kirjoittamaan koodisi siten, että samaan warppiin kuuluvat säikeet seuraavat samaa suorituspolkua
• Esimerkiksi seuraavaavat koodipätkät ovat effektiivisesti lähes yhtä kalliita:

const int local_id = get_local_id(0);
if(local_id % 2 == 0)  // Jos local_id on parillinen
    y1 = x/d;          // Jakolasku on kallis operaatio
else
    y2 = 1.0;

y1 = x/d;
y2 = 1.0;

Globaalin muistin käytöstä

• Muistipaikan osoitteen bitit määräävät välimuistin rivin, muistiohjaimen ja muistipankin (havainnekuva):

• Näytönohjaimesta riippuen muistia käytetään tehokkasti seuraten seuraavia suuntaviivoja:
  • Saman warpin/wavefrontin säikeiden tulisi jakaa mahdollisimman monta välimuistin riviä
  • Saman warpin/wavefrontin tulisi käyttää samaa muistiohjainta \(\implies\) eri warpit/wavefrontit käyttävät todennäköisesti eri muistiohjainta
  • Saman warpin/wavefrontin tulisi käyttää samaa muistipankkia \(\implies\) eri warpit/wavefrontit käyttävät todennäköisesti eri muistipankkia

• Esimerkkejä, jossa warpin/wavefrontin säikeiden muistiviittaukset valuvat usealle välimuistin riville:

• Helpoin ratkaisu on suunnitella ytimen ohjelmakoodi siten, että saman warpin/wavefrontin säikeet viittaavat vierekkäisiin N tavun sanoihin alkaen muistiosoitteesta, joka on 128 tai 256 moninkerta :

Esimerkki: summan laskeminen

1. yritys (ÄLÄ OTA MALLIA!)

• Yritetään laskea taulukon luvut yhteen GPU:lla
• Jaetaan taulukko osiin, jolloin jokainen säie voi laskea osasumman
• Lopullinen summaus voidaan tehdä CPU:lla

• OpenCL-ydin, jossa kukin säie laskee yhden osasumman ja osasummien tulokset tallennetaan res-taulukkoon (osasumman koko annetaan argumentilla m):

__kernel void sum(__global double *res, __global double *x, int n, int m) {
	const int global_id = get_global_id(0);
	
	// Lasketaan jokaiselle säikeelle osasumman rajat
	const int begin = global_id * m;
	const int end = min(n, (global_id+1) * m);
	
	// Jokainen säie laskee m peräkkäistä lukua yhteen
	double tmp = 0;
	for(int i = begin; i < end; i++)
		tmp += x[i];
	
	// Säikeet tallentavat omat osasummansa res-taulukkoon
	res[global_id] = tmp;

}

• Havainnekuva toteutuksen muistiinviittausrakenteesta (m = 8):

• Samaan warppiin kuuluvat säikeet viittaavat nyt muistiin siten, että kahden viittauksen välissä on 8 \(\times\) 8 = 64 tavua:

• Koko warpin muistihaku hajaantuu siis 16:sta 128 tavun välimuistiriville!

• Ohjelman tuloste kun taulukon koko on 43435342 ja m = 128:

$ make gpu2 && ./gpu2
g++ -std=c++11 -o gpu2 gpu2.cpp -lOpenCL
Time: 0.025615 s
Flops: 1.6957 GFlops
Sum value: 2527.55
Real value: 2527.55
Diff: 1.56433e-10

Parampi lähestymistapa

• Muokataan koodia hieman:

__kernel void sum(__global double *res, __global double *x, int n) {
	const int global_id = get_global_id(0);
	
	// Globaalin indeksiavaruuden koko (säikeiden määrä)
	const int global_size = get_global_size(0);
	
	double tmp = 0;
	
	// Säikeet aloittavat omaa indeksinumeroaan vastaavasta indeksistä ja
	// hyppäävät joka iteraatiolla global_size alkiota eteenpäin
	for(int i = global_id; i < n; i += global_size)
		tmp += x[i];
	
	// Säikeet tallentavat omat osasummansa res-taulukkoon
	res[global_id] = tmp;
}

• Yksinkertaistettu havainnekuva toteutuksen muistiinviittausrakenteesta (8 säiettä):

• Samaan warppiin kuuluvat säikeet viittaavat nyt muistiin siten, että kahden viittauksen välissä on 8 tavua:

• Koko warpin muistihaku hajaantuu siis kahdelle 128 tavun välimuistiriville!
• Huomaa, että havainnekuvaa on yksinkertaistettu!

• Testiohjelman tuloste kun taulukon koko on 43435342 ja m = 128:

$ make gpu3 && ./gpu3
g++ -std=c++11 -o gpu3 gpu3.cpp -lOpenCL
Time: 0.00473809 s
Flops: 9.16726 GFlops
Sum value: 4544.81
Real value: 4544.81
Diff: 5.63887e-10

• Toteutus on noin 5,4 kertaa nopeampi kuin ensimmäinen yritys!

Lokaalin muistin käytöstä

• Lokaali muisti on tyypillisesti jaettu 32 pankkiin siten, että peräkkäiset 32/64 bitin sanat kuuluvat peräkkäisin pankkeihin
• Jokainen pankki voi palvella yhden muistipyynnön kerrallaan (havainnekuva):

• Samaan muistipankkiin (ja eri 32/64 biti sanoihin) viittaavat muistipyynnöt käsitellään peräkkäin:

• Koko warp/wavefront joutuu odottaa jonossa olevia säikeitä!
• Tilanne, jossa saman warpin/wavefrontin säikeet viittaavat vierekkäisiin 32/64 bitin sanoihin on usein optmimaalinen lähestymistapa muistipankkien kannalta

Esimerkki: Summan laskeminen lokaalissa muistissa

• Lokaalissa muistissa oleva taulukko voidaan summata tehokkaasti käyttäen seuraavaa muistiinviittausrakennetta:

• Käydään edellinen havainnekuva hieman tarkemmin läpi
• Oletetaan, että taulukon koko on n ja, että säieryhmän koko on suurempi tai yhtäsuuri kuin n
• Taulukko jaetaan kahteen osaan: [0,m-1] ja [m,n-1]
• Luku m valitaan siten, että
  • osat ovat samankokoisia kun n on parillinen ja
  • ensimmäinen osa on yhden alkion suurempi kun n on pariton

• Seuraavaksi säikeet laskevat luvut yhteen siten, että säie, jonka lokaali indeksinumero on k laskee yhteen alkiot k ja m+k:

• Lopuksi k. säie tallentaa tuloksen muistipaikkaan k:

• Huomataan, että meillä on nyt uusi summattava taulukko, jonka koko on m \(\approx\) n/2
• Voimme toistaa saman parittaisen summauksen kunhan ensin synkronoimme säikeet!
• Muuten osa säikeistä saattaa olla vielä suorittamassa edellisen kierroksen summauksia kun osa säikeistä suorittaa jo seuraavaa iteraatiota

• Koko taulukko saadaan summattua tehokkaasti käyttäen edellä kuvattua tekniikkaa:

Globaalin ja lokaalin summauksen yhdistäminen

• Edellä kuvatut summaustekniikat voidaan yhdistää siten, että säikeet laskevat ensin osasummat globaalissa muistissa, jonka jälkeen samaan säieryhmään kuuluvat säikeet laskevat osasummansa yhteen lokaalissa muistissa:

• Testiohjelman tuloste kun taulukon koko on 43435342:

$ make gpu4 && ./gpu4
c++ -std=c++11 -o gpu4 gpu4.cpp -lOpenCL  
Time: 0.00255108 s
Flops: 17.0263 GFlops
Sum value: 1440.41
Real value: 1440.41
Diff: 1.11868e-10

• Toteutus on siis noin 1,9 kertaa nopeampi kuin edellinen toteutus!

Summan laskeminen kokonaan GPU:lla

• Summa voidaan laskea kokonaan GPU:n puolella kahden ytimen avulla:

• Vinkki: Säieryhmien koko ja määrä kannattaa valita siten, että ensimmäisen ytimen säieryhmien määrä on sama kuin toisen ytimen säieryhmän koko. Esimerkiksi:

ydin1<<<WG_SIZE_2, WG_SIZE1>>(...);
...
ydin2<<<1, WG_SIZE_2>>>(...);

• Tällöin ensimmäinen ydin laskee WG_SIZE_2 \(\times\) WG_SIZE_1 osasummaa, jotka lasketaan yhteen lokaalissa muistissa siten, että saadaan summa, jonka koko on WG_SIZE_2
• Nyt ydin2 voi laskea jäljelle jääneen summan tehokkaasti yhdellä säieryhmällä!

• Testiohjelman tuloste kun taulukon koko on 43435342:

$ make gpu5 && ./gpu5
c++ -std=c++11 -o gpu5 gpu5.cpp -lOpenCL                                  
Time: 0.00261402 s
Flops: 16.6163 GFlops
Sum value: -8903.63
Real value: -8903.63
Diff: 1.81899e-12

Yhteenveto

• Käytä paljon säikeitä ja kirjoita koodisi siten, että GPU voi hyödyntää käskytason rinnakkaisuutta
• Vältä tilanteita, joissa saman warpin/wavefrontin säikeet viittaavat muistiin hajanaisesti
• Vältä erityisesti tilanteita, joissa viittauksen välissä on 2 potenssin verran tavuja!
• Vältä tilanteita, joissa samaan warpiin/wavefrontiin kuuluvien säikeiden suorituspolut haarautuvat