Hoe theoretisch Gflops halen op fermi?

[Verwijderd]

Ik ben wat aan het klooien met OpenCL. Hiervoor heb ik een N-body simulator gemaakt, voor het testen gebruik ik N=4096. Dit algoritme heb ik zowel in OpenCL als in C++ (voor de CPU, met AVX instructies) geschreven.


CPU:
Op stackoverflow, hier en hier wordt aangegeven hoe je het maximum aantal Gflops bereikt. Op mijn sandy bridge CPU haal je dat door afwisselend een multiply en add uit te voeren. Mijn code heeft niet evenveel multiply's als add's, dus het maximum is sowieso niet te halen. Desondanks haalt mijn implementatie iets meer dan 50% van het theoretische maximum aan Gflops.

GPU:
Mijn GPU is een NVS 4200m.

Deze bestaat uit 1 SM van de fermi architectuur. Draaiend op 2x720Mhz. Het theoretisch aantal Gflops is aantal cores * kloksnelheid * 2 = 48 * 1440 * 2 = 138.24 Gflops. Diverse bronnen melden een maximum van 155Gflops. Maar het gaat daar om GPU's die op 800Mhz draaien.

Mijn kernel zit er als volgt uit:

C:

    float* xIn = hele lang array van X coordinaten;
    float* yIn = hele lang array van Y coordinaten;
    float* massIn = hele lang array van massa van de elementen

    float newdx = resulterende kracht;
    float newdy = resulterende kracht;

    for( int i = 0; i < 4096; i++ )
    {
        float xDiff = out.x - xIn[i];
        float yDiff = out.y - yIn[i];

        float x2 = xDiff * xDiff + 0.0001f; // small bias. Gratis vanwege mad (multiply-add)
        float len2 = yDiff * yDiff + x2; // mad

        float factor = massIn[i] * len2;

        newdx = xDiff * factor + newdx; // mad
        newdy = yDiff * factor + newdy; // mad

        // 16x unrolled
    }

Let op: regel 13 hoort in feite mass / sqrt( len2 )³ te zijn. Dat doen we nu even niet, dat vergemakkelijkt het rekenen.

In dit loopje worden 11 floating point berekeningen uitgevoerd. Dat gebeurt 4096² keer per frame, 2864 us per frame (volgens de profiler). We halen hierbij 64 Gflops. Dat is iets minder dan de helft van het theoretisch haalbare Gflops.

De nvidia visual profiler geeft aan dat ik compute bound ben, elke kernel 22 registers nodig heeft, en er 1024 threads naast elkaar gedraaid worden. De code zo aanpassen dat er niet over de volledige 4096 elementen geloopt wordt levert geen snelheidswinst op.


Ik heb 2 vragen:
Waarom draait deze kernel niet sneller.
Met welke magische kernel kan ik dan wél het theoretische aantal Gflops halen? De eerste stackoverflow link demonstreert hoe je dat met een CPU voor elkaar krijgt. Maar hoe doe ik dat met een GPU? Met 48k memory totaal lijkt mij cache niet het probleem (de cache L1 hit rate is ook achterlijk hoog, 99.85%)

De assembly ziet er trouwens als volgt uit, maal 16

ld.global.f32  %f29, [%r18+4]; // x
sub.f32        %f30, %f2, %f29;
ld.global.f32  %f31, [%r19+4]; // y
sub.f32        %f32, %f3, %f31;
fma.rn.f32     %f33, %f30, %f30, 0f38D1B717; // 0.0001
fma.rn.f32     %f34, %f32, %f32, %f33;
ld.global.f32  %f35, [%r20+4]; // mass
mul.f32        %f36, %f35, %f34; 
fma.rn.f32     %f37, %f30, %f36, %f27; // factor * x + dx
fma.rn.f32     %f38, %f32, %f36, %f28; // factor * y + dy

Ik neem aan dat je memory load's en multiply in één clock cycle kan doen.

[PrisonerOfPain]

Je data staat in global memory en je doet er 3 reads van - dat is sowieso je bottleneck nu. Die reads gaan absoluut niet in 1 cycle plaatsvinden (eerder 400-800 cy per read).

Je zou kunnen kijken of je je data in half-floats (.f16) kunt packen, dat scheelt je sowieso al bandwidth en verder zou je kunnen kijken of je je data in 1 read (ipv 3) binnen kunt trekken. Des noods zet je je data in shared mem (zolang het past), je leest er toch alleen maar van.

https://developer.nvidia....iciently-cuda-cc-kernelsg

[Verwijderd]

PrisonerOfPain schreef op donderdag 29 augustus 2013 @ 23:54:
Je data staat in global memory en je doet er 3 reads van - dat is sowieso je bottleneck nu. Die reads gaan absoluut niet in 1 cycle plaatsvinden (eerder 400-800 cy per read).

Zelfs als alles in L1 staat? De profiler geeft aan dat ik 99.85% L1 cache hit heb. De data is ook exact even groot als het L1.

Je zou kunnen kijken of je je data in half-floats (.f16) kunt packen, dat scheelt je sowieso al bandwidth en verder zou je kunnen kijken of je je data in 1 read (ipv 3) binnen kunt trekken. Des noods zet je je data in shared mem (zolang het past), je leest er toch alleen maar van.

Als ik de loads omtover naar half floats, daalt de framerate licht. Volgens mij is de memory access echt geen bottleneck. CPU-z geeft aan dat ik nul procent van de memory controller gebruik. De pofiler geeft aan dat ik 0.17GB/s memory bandwidth gebruik.

Het reorderen van de data op de manier die jij vermeld levert een negatief effect op de performance (ongeveer 8%). De code vectorizen (float4/8, vload4/8) zodat er eventueel efficiënter geload kan worden levert eveneens een performance regressie (vanwege de hogere register load?)

Trouwens, een aantal van mijn kernel argumenten waren nog __global in plaats van __constant. Dat laatste levert een aanzienlijke performancewinst op. OpenCL staat het met mijn programma vreemd genoeg niet toe om non-primitive arrays als const te markeren. Nochtans kom ik zelfs met deze optimalisaties niet boven de 50% van de theoretische flops uit.

[Zoijar]

Het leek mij ook als ik de asm boven zie dat je nooit je loads kan latency hiden met zo weinig compute tegenover zoveel load. *de* bottleneck op fermi is memory access. Maar blijkbaar geeft hij aan van wel... ik zou haast vragen wat maakt cuda ervan.

[Verwijderd]

Zoijar schreef op vrijdag 30 augustus 2013 @ 08:19:
Het leek mij ook als ik de asm boven zie dat je nooit je loads kan latency hiden met zo weinig compute tegenover zoveel load.

Volgens de profiler scheduet hij 1024 threads naast elkaar. In principe is het zo dat alle 1024 threads precies hetzelfde memory blokje proberen te accessen. Misschien dat er daarom zo weinig memory load is?

[PrisonerOfPain]

Verwijderd schreef op vrijdag 30 augustus 2013 @ 01:45:
[...]

Zelfs als alles in L1 staat? De profiler geeft aan dat ik 99.85% L1 cache hit heb. De data is ook exact even groot als het L1.

L1 is dezelfde chip als je shared memory, maar ook dat heb je niet instant; volgens deze thread heb je alsnog een 88cy latency op je L1. Dus ondanks dat je snellere access hebt ben je nog steeds keihard mem-bound.

Verder zou je hier nog kunnen kijken naar verdere tip over hoe je je data acess zou kunnen verbeteren.

Darkstone schreef op vrijdag 30 augustus 2013 @ 11:15:
[...]

Volgens de profiler scheduet hij 1024 threads naast elkaar. In principe is het zo dat alle 1024 threads precies hetzelfde memory blokje proberen te accessen. Misschien dat er daarom zo weinig memory load is?

Kijk eens of je geen bank conflicts op je shared mem hebt, dan

[Verwijderd]

Volgens mij hebben jullie gelijk dat ik memory bottlenecked ben. Als ik een paar extra MAD's in mijn loop plak haal ik 80% van de theoretische flops. Ik had verwacht dat de latency hiding beter zou werken, maar schijnbaar is dat nog steeds een bottleneck. Ik ga de data nog een paar keer reorderen om te kijken of er nog iets optimaliers te vinden is dan wat ik nu gebruik.

Verder heb ik volgens mij geen bank conflicts, omdat elke thread van hetzelfde address in de bank probeert te lezen. Als ik de loop niet bij 0 begin maar bij een veelvoud van 2 op basis van het thread id, daalt de framerate sterk als dat veelvoud klein is. De visual profiler laat helaas geen bank conflicts zien. Waarschijnlijk omdat mijn GPU redelijk antiek is..

Waarbij de Y as de framerate is en de X as de macht van 2. Bij X <= 8 leest hij een gedeelte van de array out of bounds, maar dat heeft geen effect op de performance volgens een snelle test.

De loop ziet er dan zo uit:

C:

const int id = get_global_id( 0 );

int endID = id & ~1023; // of andere macht van 2 - 1
int i = endID;
do {
  
    ... doe iets
    i++;
    16x unrolld

  i = i%4096;
} while ( i != endID );

Je ziet dat de efficiëntie een procent of 2 a 3 beter is wanneer X iets kleiner is dan het aantal threads, zodat er van 2, 4 of 8 verschillende adressen gelezen wordt in plaats van dat alle threads hetzelfde adres lezen. Niet erg significant.

[Zoijar]

Verwijderd schreef op vrijdag 30 augustus 2013 @ 13:16:
Verder heb ik volgens mij geen bank conflicts, omdat elke thread van hetzelfde address in de bank probeert te lezen.

Om coalseced reads goed te laten gaan moet elke thread juist van een consecutive address lezen op hetzelfde moment. Threads leven niet als normale cpu threads: gpu threads draaien volledig in lockstep (per warp), alsof het een hele brede simd architectuur is (ala SSE). Als er wordt ontdekt dat bij een load thread 1 van n leest, threads 2 van n+1, etc, dan wordt er een enkele memory transactie uitgevoerd (zeg maar 1 shared transactie die 32 waarden ophaalt, voor elke thread 1, in plaats van 32 losse transacties voor elke thread van 1 waarde) Het gaat niet eens zozeer om throughput oid, het gaat om transacties die zo duur zijn, of je nou 4 of 256 (whatever max) bytes leest.

Daarnaast heb je te maken met occupancy: zodra er een transactie nodig is scheduled hij een andere runable warp. Er moeten genoeg warp actief/runable/in-flight zijn om de hele latency te kunnen hiden. Dus a) minder latency maken door coalseced reads en b) de latency die je krijgt voldoende overlappen. Dat is eigenlijk het hele geheim van high-performance gpgpu.


Maar ik ben een beetje roestig wat betreft gpgpu programmeren, dus misschien begrijp ik je verkeerd

[ Voor 5% gewijzigd door Zoijar op 30-08-2013 18:07 ]

[Verwijderd]

Coalesced reads is een term die wordt gebruikt bij communicatie naar het global memory toe. Bank conflicts is een term die wordt gebruikt bij shared (on chip, L1 of constant) memory. Ik ben er alleen nog niet helemaal achter wat de compiler nu vertaalt naar shared memory of global memory.

Aan het begin van de shader wordt er éénmaal van het global memory gelezen om de huidige body (struct van 5 floats) op te halen. aan het einde van de shader worden deze 5 floats weer weggeschreven. Ik denk, maar weet niet zeker, dat de kosten van deze memory access in het niet vallen naast die gigantische loop waar per shader ~40k floating point berekeningen worden uitgevoerd.

Ik neem aan dat de bottleneck het lezen van xIn, yIn en massIn is. Dat levert geen bank conflicts op, omdat elke thread in de warp van dezelfde 'i' leest. (zie loop in topicstart). Ik neem dan even aan dat een thread uit een warp niet op een andere thread voor mag lopen.

Ik heb nog wat geprobeerd te knutselen met de memory layout. Losse arrays voor X, Y en mass werken het best. Een memory layouy als:

C++:

Body {
  float x;
  float y;
  float mass;
}

In plaats van de losse arrays is ~10% trager. Als ik padding toe voeg wordt het nog veel trager.


Edit: bij nader inzien bestaat we wel een bank conflict getalletje in de visual profiler. Ik weet niet of hij correcte resultaten weergeeft, maar hij zegt dat er geen enkele bank conflict is tijdens de run. Ik vind de visual profiler wat lastig om mee te werken. De resultaten zijn nogal abstract..

[PrisonerOfPain]

Verwijderd schreef op vrijdag 30 augustus 2013 @ 19:15:
Coalesced reads is een term die wordt gebruikt bij communicatie naar het global memory toe. Bank conflicts is een term die wordt gebruikt bij shared (on chip, L1 of constant) memory. Ik ben er alleen nog niet helemaal achter wat de compiler nu vertaalt naar shared memory of global memory.

In princiepe beheer je je shared memory zelf, je kunt de configuratie zelf instellen, of 16KB L1 of 48KB shared mem of andersom. Je L1 word automatisch gepopulate door reads naar global memory te doen terwijl jij de baas bent over je shared mem.

Post eens een screenshot van je profiler / of een set resultaten? Nu zitten we maar een beetje naar halve code-snippets te kijken terwijl bij optimization meestal de hele code + de profiler het belangrijkste zijn.

[Verwijderd]

Ik had nog wat geprobeerd te profile, maar dat wil nog niet echt lukken. Met name door een gebrek aan goede tools. De visual profiler 4.0 geeft na het profilen een error, en dropt vervolgens 3800 van de 3810 (ofzo) resultaten.. Verder zijn de cijfers bij de resultaten 3 maal zo hoog als ze eerder waren. Daar klopt geen kant van.

Versie 4.1 wil het op mijn computer niet doen. Vanaf 4.2 heeft nvidia de support voor OpenCL er uit geflikkerd.

Aan de kant van AMD hebben we CodeXL. Die zorgt direct voor een bluescreen als ik een profile wil starten. Dat schiet ook niet op.

Daarom gaat dit project in de ijskast. Ik wil graag goede tools om de bottlenecks in mijn applicatie te vinden, maar die bestaan blijkbaar niet.