Grafikkort till parallella beräkningar

Teknik och samhälle Datavetenskap

Examensarbete

15 högskolepoäng, grundnivå

Grafikkort till parallella beräkningar

Graphics cards for parallel computations

Författare: Sani Music

Examen: Högskoleingenjörsexamen 180 hp Handledare: Mia Persson

Huvudämne: Datavetenskap Examinator: Ivan Kruzela

Program: Data och telekom

Resumé

Den här studien beskriver hur grafikkort kan användas på en bredare front än multimedia. Arbetet förklarar och diskuterar huvudsakliga alternativ som finns till att använda grafikkort till generella operationer i dagsläget. Inom denna stu-die används Nvidias CUDA arkitektur. Stustu-dien beskriver hur grafikkort används till egna operationer rent praktiskt ur perspektivet att vi redan kan programme-ra i högnivåspråk och har grundläggande kunskap om hur en dator fungeprogramme-rar. Vi använder s.k. accelererade bibliotek på grafikkortet (THRUST och CUBLAS) för att uppnå målet som är utveckling av programvara och prestandatest. Resulta-tet är program som använder GPU:n till generella och prestandatest av dessa, för lösning av olika problem (matrismultiplikation, sortering, binärsökning och vektor-inventering) där grafikkortet jämförs med processorn seriellt och paral-lellt. Resultat visar att grafikkortet exekverar upp till ungefär 50 gånger snab-bare (tidsmässigt) kod jämfört med seriella program på processorn.

Abstract

This study describes how we can use graphics cards for general purpose compu-ting which differs from the most usual field where graphics cards are used, mul-timedia. The study describes and discusses present day alternatives for using graphic cards for general operations. In this study we use and describe Nvidia CUDA architecture. The study describes how we can use graphic cards for ge-neral operations from the point of view that we have programming knowledge in some high-level programming language and knowledge of how a computer works. We use accelerated libraries (THRUST and CUBLAS) to achieve our goals on the graphics card, which are software development and benchmarking. The results are programs countering certain problems (matrix multiplication, sor-ting, binary search, vector inverting) and the execution time and speedup for these programs. The graphics card is compared to the processor in serial and the processor in parallel. Results show a speedup of up to approximatly 50 times compared to serial implementations on the processor.

Centrala nyckelord: Nvidia CUDA, THRUST, CUBLAS, Eigen, OpenMP,

accelererade bibliotek, prestandatest, GPU, CPU, vektor inventering, sortering, binärsökning, matrismultiplikation

Innehåll

3.2 HyperThreading och speedUP . . . 7

3.3 Grafikkort . . . 8

3.4 GPU . . . 8

3.5 Grafikkortets minne . . . 8

3.6 Grafikkortet till generella beräkningar . . . 8

3.7 Standard template library . . . 9

3.8 Binärsökning . . . 9

3.9 Sortering . . . 9

3.10 Vektorinventering . . . 10

3.11 Matrismultiplikation . . . 10

3.12 Eigen . . . 10

4 Fördjupning i Nvidia CUDA 11 4.1 CUDA C/C++ . . . 11

4.2 Minneshantering i CUDA . . . 11

4.3 Kernel i CUDA . . . 11

4.4 Parallellism i CUDA . . . 12

4.4.1 Synkrona och asynkrona CUDA exekveringar . . . 12

4.5 Accelererade bibliotek och THRUST . . . 12

4.5.1 THRUST . . . 12

4.5.2 CUBLAS . . . 13

5 Experiment 14 5.1 Vektorgeneratorn . . . 15

5.2 Struktur på vektorhanterings program . . . 16

5.3 Parallellt på CPU med OpenMP . . . 19

5.4 Sortering med THRUST och STL . . . 19

5.5 Binärsökning med THRUST och STL . . . 20

5.7 Matrismultiplikation av kvadratiska matriser . . . 21 5.7.1 Matrismultiplikation på CPU . . . 21 5.7.2 Matrismultiplikation på GPU med CUBLAS . . . 22

6 Diskussion 24

6.1 Analys av vektor program . . . 24 6.2 Analys av matrismultiplikation . . . 27

7 Slutsats 30

7.1 Vidare studier . . . 30

A Mätdata vektorhanterings program 35

1

Inledning

Spel och annan form av multimedia driver utveckling av grafikkort fort de senaste åren. Grafikkorten tillgängliga på marknaden i nutid är mer kraftfulla och utbudet är större jämfört med förr [21]. Ifall vi jämför grafiken i spel idag med 10 år gamla spel ser vi en visuell förbättring, grafiken (flöde av bilder) skapas på grafikkortet [11]. Om vi omdirigerar kapaciteten som går åt att skapa dessa bilder till andra om-råden än multimedia får vi någon form av prestanda. Det är denna prestanda och tillväga som diskuteras och analyseras i uppsatsen. Detta är intressant eftersom det undersöker sätt att ge ökad prestanda med hårdvara ofta befintlig i stationära och bärbara datorer.

Grafikkort är parallella enheter som utvecklats till ”many-core” (Nvidia GT550M har 144 kärnor) system. Algoritmer som går att tillämpa på grafikkortet är parallel-la algoritmer. Visa av dessa algoritmer utveckparallel-lades under 80 talet för att användas på (då) dyra svåråtkomliga (då) superdatorer. Dessa algoritmer går att implemen-tera på grafikkort idag eftersom dessa är hårdvarumässigt byggda för many-core parallellism. [21]

Grafikkort går att använda till generella beräkningar (annat än multimedia). Fors-kare och utvecklare inser speciellt med lansering av Nvidia Geforce 3 kretsen (år 2001) ”råkraften” grafikkort har i generell beräknings kapacitet. Men att använda denna råkraft till annat än grafiska uppgifter är en då tidskrävande process. Ef-tersom det endast är möjligt genom att använda grafiska APIs (OpenGL, DirectX) och ”lura” grafikkortet att utföra generella beräkningar. Detta förändras när Nvidia år 2006 (tillsammans med Geforce 8 serien) lanserar sin CUDA arkitektur, speciellt konstruerad att tillåta grafikkort att utföra generella beräkningar. [18]

Nvidia menar att det ska gå att få betydligt snabbare programvara vid använd-ning av CUDA arkitekturen jämfört med seriella implementationer, tresiffriga tal används t.ex. 2600 gånger snabbare [42]. Wynters visar i artikeln ”Parallell proces-sing on Nvidia graphics cards” [53] parallellism rika fall där grafikkort med Nvidia CUDA teknologi presterar upp till 2412 gånger snabbare än motsvarig seriell im-plementation på processorn.

1.1 Forskningsfråga

Har tiden kommit då det finns en aritmetisk superdator i de flesta hem och i så fall hur lätt är det att använda den? Hur mycket tidsmässigt snabbare är paral-lellt exekverbar programvara på ett Nvidia GT550M grafikkort jämfört med en Intel 2670QM i7 processorn med hänsyn till parallellt och seriellt exekverade program på processorn?

1.2 Syfte

Denna studie syftar till att ge utvecklare och andra intressenter en introduktion, vägledning och mätdata i processen att utnyttja GPU-kapacitet i datorer. Samt för-klara vilka alternativ som finns i fältet. Studien syftar att presentera program och prestandamätdata av dessa som är baserad på CUDA teknologi. Mätdata som pro-duceras och tillvägagångsätt för denna är centrum för studien och det är resultatet av dessa som ska hjälpa intressenter i eventuellt val av implementering av CUDA i programvara.

1.3 Målgrupp

Delar av denna uppsats är främst till läsare som redan har programmerings kun-skap i högnivåspråk helst C, C++ (p.g.a. b.la. minneshantering) dessutom god da-torvana och förståelse för grundbegrepp inom datorer. Läsare som inte har dessa kunskaper kan studera diskussions avsnittet utan att förstå studien i detalj.

1.4 Avgränsningar

Tabellerna här visar vilken mjukvara respektive hårdvara som används i studien. Samtliga experiment i denna uppsats använder denna hårdvara och mjukvara.

Tabell 1: Hårdvaruspecifikation i denna studie

Tillverkare Typ Namn Relevant Information

Intel Processor Core i7 2670QM 32nm (2nd gen) 2.2GHZ, 6MB Cachemin-ne, 4 kärnor

Nvidia Grafikkort GT 550M (40nm) 144 CUDA cores (1.48GHZ GPU), 1GB GDRAM

Samsung Ram RAM DDR-3 6144 MB 1.33 GHZ DDR3

SDRAM-minne [37], [45], [16].

Tabell 2: Mjukvaraspecifikation i denna studie

Leverantör Typ Namn Relevant Information

Microsoft Operativsystem Windows 7 SP1

Microsoft Kompilator Visual Studio C++ Ver 2010 Ultimate edition

OpenMP ARB API Open MP Integrerad i kompilator

Nvidia Arkitektur CUDA V4.1

Nvidia Bibliotek CUBLAS V4.1

THRUST Bibliotek THRUST V1.51 (Inkl i CUDA)

Eigen Bibliotek Eigen V3.05

[34], [13], [9], [40].

andra teknologier diskuteras kort men behandlas ej i denna uppsats p.g.a. uppsat-sens omfattning i tid.

2

Metod

Vi diskuterar några möjliga empiriska studier som tas upp av Heiko Koziolek i ”The Role of Experimentation in Software engineering”[24]. Kontrollerade experiment är metoden med högst pålitlighet eftersom den tillåter mest kontroll över variabler som kan påverka resultatet förutom den oberoende variabeln. En mindre strikt metod är fallstudien där ett specialfall analyseras, initiala förhållande kan analyseras men slutsatsen är inte generaliserbar. Metastudier analyserar andra studier för att se-dan jämföra resultaten.

Experiment är byggda runt en ”mini” verklighet vilket innebär att det handlar om en förenklad bild av verkligheten där problemen är isolerade (Många variabler ab-straheras bort). Fördelen att använda experiment är återupprepningsbarheten och att vi har stor kontroll över variabler som påverkar studien. [4]

2.1 Metodval

Vi väljer en kvantitativ empiriska studie med kontrollerade experiment där mätdata samlas och analyseras. Experimenten och resultatet av dem är centralt för studien och det viktigaste. Teorin som används finns i kapitel 3, 4 och motiverar detaljerat vilka val som tas praktiskt och varför just dessa anses vara bäst. Kapitel 5 visar de-taljerat experimenten, hur dessa utförs (mätdata finns som appendix). Experiment kapitlet diskuterar val som tas i detalj med experimenten och motiverar varför des-sa val tas. Resultatet vi strävar efter är av ordes-sak-verkan typ.

Vi skriver egna CUDA applikationer som vi sedan jämför med likvärdiga egenskriv-na C++ program i experiment och observerar tiden (mätdata) som krävs för olika subrutiner att exekvera. Det skulle även gå att använda existerande program och studier för att uppnå syftet (utan hårdvaru avgränsningen). Men detta sätt är inte lika intressant då det inte medför någon ny mätdata till området.

Genom experiment ser vi en praktisk implementation av teorin och hårdvaran vi använder. Vi skriver inte program ämnade åt absoluta slutanvändare då denna typ av programvara kräver resurser utanför vår tidsram d.v.s. accepterar vi att våra program inte är absolut felfria.

2.1.1 Experimentkonstruktion

Det finns några olika sätt att uppnå syftet praktiskt med kontrollerade experiment. Ett sätt är att skriva egna algoritmer på processorn och grafikkortet och sedan jäm-föra exekverings tid av dessa. Detta är det svåraste sättet och det som kräver mest djup förståelse för b.la. hårdvara och mjukvara. Syftet skulle gå att uppfylla på det-ta sätt men det är inte ”best practice” eftersom det inte visar hänsyn till teorin (se kapitel 3,4) från b.la Nvidia och tidigare studier speciellt Langdons ”debuging cuda” [25]. Fördelen med detta sätt är att vi kan skapa nya algoritmer kanske bättre än de som existerar i specialfall.

I denna studie väljs fyra problem att behandla på grafikkort och processorn. Des-sa är binärsökning, vektor inventering, sortering, matrismultiplikation. Samtliga av dessa fyra problem är återkommande i datavetenskap. Vi väljer i denna uppsats att arbeta med existerande bibliotek och algoritmer (CUDA och CUBLAS på grafik-kortet, STL och Eigen på processorn) i full möjlighet då detta ger kvalitativ kod att utföra experiment med och bör därför ge mer relevant mätdata för forskningsområ-det och intressenter.

2.2 Mätdata

Mätdata byggs runt tid som är oberoende av resten av experiment och går att mäta med instrument helt oberoende av hårdvaran. Prestandatest är upprepningsbara fast med hänsyn till omständigheter (OS bakgrundsaktiviteter) kan avvikelse ske. Mätdata sparas i tabeller bifogade till uppsatsen. Ur mätdata kan vi skapa nya ta-beller som ger översiktlig information. Mätdata innehåller alltid ett unikt s.k. ”Test ID” som skiljer exekveringen av det specifika experimentet från resten. Vi tar i be-aktning eventuella felkällor i vår mätdata och försökt minimera dess påverkan. Den främsta felkällan som kan störa vår data är att CPU:n är belastad under fram-tagning av mätdata respektive GPU:n men detta undviks genom att överbevaka systemresurser och inte använda datorn till något annat än experiment vid datain-samling. Vidare så är en annan störning att när vi kallar på GPU:n så är kallelsen asynkron dvs återlämnar GPU exekveringen kontroll till CPU omedelbart efter ex-ekvering. För att kunna ta verklig tidstagning följer vi rekommendationer i "CUDA best practices” och gör tidtagning synkroniserat (detaljer finns i teori avsnitt).

2.3 Litteratur

Vi använder relevant litteratur i form av internetsidor, böcker, material från Nvidia och vetenskapligt material för att utveckla teori om grafikkort till generella beräk-ningar. Speciellt används två böcker ”Programming massivly parallell programs” [21] av Kirk, Hwu och ”introduction by Example to CUDA” [18] av Kandrot och Sanders. Från Nvidia är ”Cuda C Best Practice Guide” [35], ”Cuda C Programming Guide” [39], ”Cublas Documentation V4.1” [41] intressant material.

2.4 Metodkritik

Experiment ger hög validitet men kan anses vara en förenklad världsbild som inte ger full i insikt och hänsyn till omgivningen, experiment kan vara dyra att utfö-ra. Ibland missanvänds statistik av mjukvaru experiment. Experimenten kommer endast ge en djup förståelse för den hårdvara som definierats i avgränsningen ef-tersom det är på denna som experimenten kommer ske. Vi visar (Kapitel 5) i detalj hur experimenten utförs (vilken kod som används), vilket innebär att samma expe-riment skulle gå enkelt att utföra på annan hårdvara.

3

Teknisk bakgrund

Det är processorn som utför instruktioner i en dator och vi associerar den ofta till hjärnan i datorn. Här refererar vi till processorer som CPU (Central Processing Unit). CPU:n är designat för kontroll av data flöde och det är processorn som i ”nor-mala” fall utför beräkningar på datorn.

Cacheminne ligger fysiskt på processor chippet, cacheminnet är relativt litet men snabbt att skriva och läsa till. RAM (Random Access Memory) minnet lagrar kör data och är mycket större än cacheminnet. RAM är samman kopplat till processorn via något gränssnitt i detta fall Intel’s ”Quick path” [1]. Processordelen som ut-för aritmetiska instruktioner och beräkningar ut-förkorttas ALU (Aithmetichal Logical Unit). Intel i7 Processorn lägger mer vikt på cacheminnet än på ALU. Anledningen till att detta är bl.a. att det minskar trafik på ”quickpath” och efterdröjning mellan CPU och RAM. [27]

3.1 Parallellism

Vi associerar antagligen parallellt med någonting enkelt som t.ex. två räta linjer som inte korsar varandra, på sätt och viss är två parallella program två linjer som inte korsar varandra. Dessa två linjer är tillåtna att kommunicera, och även dela da-ta under förutsättning att de inte korsar varandra. Detda-ta känns naturligt eftersom t.ex. två program inte kan samtidigt finnas på identiska adresser i RAM minnet och vara olika.

Amdhals lag beräknar hur mycket snabbare en applikation blir ifall källkoden pa-rallelliseras.

Snabbarekod=

1

F + (1−F )_N (1)

Andel av koden som inte går att parallellisera betecknas som F . Ett illustrerande exempel på detta är i fallet vi uppnår 50 procent parallellism då tilldelas F värdet 0.5. Variabeln N beskriver antalet kärnor som används. Vi måste ha en kärna för att exekvera applikationen. Om det finns 10 parallella kärnor tilldelas N värdet 10. Med dessa variabeltilldelningar så erhåller vi följande värden.

Snabbarekod=

1

0.5 + 1−0.5₁₀ = 1.8181...

En applikation som är nästan två gånger snabbare p.g.a. parallellism kan verka som mycket men vi inser snabbt att med ett lägre F och större N kan det bli be-tydligt mer. Testar vi olika värden med Amdhals lag ser vi snabbt att antalet trådar och andel parallelliserad kod samverkar och att det är oerhört viktigt att ha en stor

andel av koden parallelliserad (litet värde på F ). Följande bild visar en kontroll i ett program som skulle vara helt meningslös att köra i parellallism.

Figur 1: Kontroll av födelsedatum

Det finns ingen mening i att försöka uppnå parallellism vid inmatning av t.ex. födelsedatum från en användare. Även Ifall vi uppnår en parallellism på 100 pro-cent så är det meningslöst. Problemet är inte av den typ vi kan bemöta med parallell databehandling eftersom det inte finns data att behandla parallellt.[2]

Denna studie använder parallellism på CPU:n med OpenMP API:et som redan finns tillgängligt i kompilatorn, men vi måste aktivera användning av OpenMP i kompi-latorn [44]. I denna studie refereras till parallell exekvering på CPU:n som ”CPU/P” och seriell som ”CPU/S”.

Trådsynkroniseringen i denna studie är av typen ”nowait” vilket innebär att trå-dar exekverar samtidigt, detta konfigureras genom att inte ge OpenMP s.k. ”direktiv klamrar” [31].

3.2 HyperThreading och speedUP

Intel processorer levereras med en teknologi som heter ”HyperThreading” (HT). I vårt fall har vi inte 8 processor kärnor trots att operativsystemet tror att vi har det. Denna hårdvara har 4 riktiga och 4 emulerade kärnor för prestanda d.v.s. finns 2 emulerade processor trådar per varje en riktig kärna. [15]

Men ur operativsystemets perspektiv har vi 8 stycken kärnor. Detta är betydelse-fullt så att vi vet vilka variabelvärden att använda i Amdahls formel när vi försöker förutse hur mycket snabbare mjukvara blir på processorn. Om vi lyckas uppnå 100 procent parallellism kommer programmen inte bli 8 gånger snabbare men de kan bli snabbare än 4 gånger beroende på hur mycket HT som kan tillämpas. Proces-sorn som används i studien har ytterligare en teknologi vi måste förstå som heter ”SpeedUP” och möjliggör att processorn vid intensiva arbete kan öka frekvensen upp t.o.m. 3.10GHZ från de ursprungliga 2.20GHZ. [14] [16]

3.3 Grafikkort

Det är grafikkortets uppgift att skapa och visa bilder på skärmen. Ett modernt grafikkort kopplas till datorns moderkort genom gränssnittet ”PCI-E 2” (Peripheral Component Interconnect Express 2) [45]. Det finns grafikkort som använder andra gränssnitt men här har vi ett separerat grafikkort kopplat via PCI-E 2 [37].

3.4 GPU

Här refererar vi till grafikkortets processor ”Graphical Processing Unit” som GPU. Till skillnad från CPU är GPU dedikerat åt parallella databeräkningar. En GPU be-står av många ”multi cores”, kärnor på svenska som exekverar parallellt kod. Varje multicore har bland annat 8 ”CUDA cores” vi har 18 multicores med denna hårdva-ra eftersom denna hårdvahårdva-ra har 144 CUDA cores. GPU:n har sitt eget cacheminne på chippet. [39]

3.5 Grafikkortets minne

Ett grafikkort har sitt eget RAM-minne separerat från RAM-minnet på moderkortet. RAM-minnet lagrar data av olika typer som t.ex. färdigställda bilder, positioner på skärmen etc. De flesta grafikkort som lanseras idag har över 1024 Megabyte RAM minne [23]. På GT550M kortet är gränssnittet mellan GPU och grafikkortets RAM 28.8GB/sec [37].

Bandbredden på PCIE 2 interfacet ligger på 8GB/s [45], av bl.a. denna anledning är det eftersträvat att hålla data så länge som möjligt på grafikkortet och undvika kommunikation mellan CPU:n och GPU:n även om det inte finns någon kortsiktig prestanda-vinst på att utföra subrutiner på GPU:n . Det finns fler typer av minne på ett GT550M kort, cacheminne är det minnet som ligger på GPU:n och är extremt snabbt, cacheminne är dock extremt lite tillgängligt precis som på CPU:n. [35] [21]

3.6 Grafikkortet till generella beräkningar

Det finns olika alternativ till att programmera grafikkort att utföra mer generella uppgifter än multimedia, dessa uppgifter kan vara helt orelaterade till grafik. De tre dominerande alternativen är Nvidias CUDA, OpenCL och Microsofts Directcompute

[20]. Här väljer vi Nvidia CUDA teknologin p.g.a. allt material som erbjuds av Nvi-dia och lätt tillgängligheten av GPU som stödjer CUDA samt hårdvaran på systemet som används. Nvidia har inte exkluderat användning av OpenCL och Directcompute på deras hårdvara [19]. CUDA förkortat för ”Compute Unified Device Architecture” och är Nvidias modell till att använda GPU:s och integrerade komponenter till icke grafiska operationer. CUDA kräver ingen licens och kostar ingenting. CUDA började lanseras år 2007 tillsammans med Nvidia 8 serien och har levererats med över 100 miljoner GPU enheter [21]. Ett externt grafikkort med CUDA kan idag kosta mellan några hundra kronor och flera tusen enligt priser på webbsidan www.komplett.se [22]. Förutom ett Nvidia grafikkort med CUDA stöd behövs inga andra specifika hårdvara komponenter för att använda CUDA. Priset och Nvidias marknads ande-lar gör att CUDA teknologi är lättillgänglig. CUDA C/C++ är en utvidgad version av programmerings språket C och är standard språket som används till att program-mera GPU enheter som stödjer CUDA. [32]. Det finns även tredjeparts tillägg som möjliggör att Cuda används med programmerings språk så som JAVA SE, C# [17], i denna studie används enbart CUDA C/C++.

3.7 Standard template library

STL (Standard Template Library) är ett bibliotek som tillhör C++ standardbiblioteket och erbjuder många algoritmer, datastrukturer, klasser etc. som ofta återanvänds. STL är ett generiskt bibliotek vilket innebär att det är skrivet på ett sätt att vi kan använda samma kod med olika datatyper. [51]

3.8 Binärsökning

Binär sökning är ett sätt att söka igenom en sorterad lista genom att halvera listan vid varje ”försök”. Om vi har 1 miljon sorterade element i en lista och ska hitta just ett element kommer listan att halveras vid varje försök. Exempelvis om talet 77 ska hittas hade sökningen börjat på element 500 000 och undersökt om elementet är större än eller mindre än 77. Om talet är större jämförs 750 000 annars 250 000 etc. Komplexiteten för en sådan sökning är O(logN ) där den sökta listans längd är N. Vanligen är komplexiteten i vårt fall 6 sökningar men i bästa fall O(1) om talet hittas just i mitten. Det finns en funktion i STL för binärsökning på vektorer. [5], [30]

3.9 Sortering

Det finns olika algoritmer i datavetenskap att sortera data med, i olika situationer kan vissa algoritmer vara bättre val än andra då de erbjuder t.ex. stabilitet eller snabbhet. I denna studie används sorteringsfunktioner från STL (sort respektive stable_sort). Sort har en vanlig komplexitet på N logN och i värsta fall N2_{, stable_sort}

3.10 Vektorinventering

Vektor inventering betyder att vektorn vänder på sina element eller en del av dem. Om vektor A innehåller elementen 1, 3, 4, 5, 7 skulle vektor A inventerat bli 7, 5, 4, 3, 1. Det finns en funktion för detta i STL och den har en linjär komplexitet O(N ). [8], [47]

3.11 Matrismultiplikation

Matriser används mycket i datavetenskap t.ex. i bildbehandling och beräkningar i rummet. I denna studie behandlas tvådimensionella matriser de multipliceras med varandra på CPU:n och GPU:n och resultaten jämförs. Om vi har två matriser A och B och vill multiplicera dem med varandra måste vi kontrollera att A har lika många kolumner som B har rader. Om C är resultatet av AB så gäller för C följande

Ci,j = Ai,1B1,j+ Ai,2B2,j+ ... + Ai,nBn,j = n

X

r=1

Ai,rBr,j (2)

Där i anger rader och j kolumner. Ett exempel: A =1 2 3 4  , B =0 5 2 6  C = AB C =((1 · 0) + (2 · 2)) ((1 · 5) + (2 · 6)) ((3 · 0) + (4 · 2)) ((3 · 5) + (4 · 6))  =4 17 8 39  [48]

Dessa matriser är ”dense” matriser eftersom de innehåller data i majoriteten av elementen. Alternativet är ”sparse” matriser som innehåller majoriteten nollor [10]. I denna studie behandlar vi endast ”dense” matriser. Komplexiteten för matrismul-tiplikation av nxn storlek är O(N3) [52].

3.12 Eigen

Matrismultiplikation verkar som en enkel uppgift att implementera i t.ex. C/C++. Allt vi behöver är tre iterationer, problemet med denna retorik är att matrismulti-plikation (likt mycket annat) kräver en hel del hänsyn och djup förståelse till minne för att vara optimerat på hårdvaran som används. Det finns detaljerad text om själ-va optimeringen av Drepper [50]. Eftersom de matrisoperationer som behandlas i experiment här är återkommande och vanliga använder vi existerande, optimerade bibliotek. Vi väljer att använda Eigen biblioteket till CPU program (Publicerat med GNU free document license) p.g.a. erhållen prestanda mätning av matris-matris multiplikation på projektets sida. Eigen är enkelt att implementera och använda, för att implementera Eigen inkluderas relevanta filer (C++). [9]

4

Fördjupning i Nvidia CUDA

Vi går kort igenom CUDA i detta introducerande avsnitt, för detaljer rekommen-deras litteraturen som används här. För att kunna utveckla i CUDA hämtas och installeras först nödvändiga komponenter (SDK, drivrutiner och toolkit) som finns tillgängliga kostnadsfritt på Nvidias hemsida [42]. Det medföljer vissa filer som hjäl-per oss att snabbt komma igång med CUDA. [39],[35],[36]

I denna studie används Visual studio 10 Ultimate Edition som kompilator. För att separera CPU och GPU kod inbäddas en såkallad ”Nvidia C Compiler” (nvcc) i ut-vecklings miljön. Källkod som används av GPU:n levereras i ”PTX” format. Kompi-latorn vet om att den ska använda CUDA p.g.a. fil ändelsen till CUDA filer är ”*.cu”. CUDA är designat att tillåta skalbarhet utan att utvecklare oroar sig för hur samma kod skall exekvera på olika GPU. Detta fungerar därför att Nvidia C kompilatorn levererar kod i PTX som kompileras vid kör tid. [33]

4.1 CUDA C/C++

Nvidia utvecklar CUDA C/C++ (Oftast kallat CUDA C) som ett utbrett språk till pro-grammering språket C med ”minimalt” nödvändig utveckling. Syntaxen är samma som C++ och går att (som med vanlig C kod) blanda med C++ källkod. [18]

4.2 Minneshantering i CUDA

Vi måste med CUDA allokera respektive deallokera grafikkorts-minne men även skicka minne fram och tillbaka mellan CPU:n och GPU:n eftersom grafikkort inte kan hantera sitt eget minne. Vi allokerar och deallokerar minne med specifika in-struktioner till CUDA C/C++ i denna studie. [35]

Enligt bl.a. tidigare studie av Langdon [26] är det viktigt att undvika data-kommunikation till och från GPU:n p.g.a att detta är en tidskrävande operation. Ifall problem som går att bemöta med GPU:n (parallella problem) är databeroende (filer på t.ex. hård-disken) kan de behandlas relativt snabbt p.g.a. GPU:ns parallella struktur. Men det blir svårt att konstant skicka in data ifall problemet inte är datarikt.

4.3 Kernel i CUDA

I detta sammanhang är funktioner som anropas från CPU men exekveras på GPU benämnda ”Kernels”. Endast en Kernel exekveras åt gången i experiment här. Anrop till Kernel är asynkrona och återlämnar kontroll till CPU:n efter några mikrosekun-der detta är viktigt att förstå vid tidtagning. En Kernel börjar först exekvera när föregående CUDA operationer är färdiga (med vårt system). Vid en exekvering av en Kernel i CUDA C definieras antal trådar, blocks, grids där dessa tre är relaterade till problemets ”storlek”, vid användning av bibliotek är detta inte nödvändigt. [39]

4.4 Parallellism i CUDA

CUDA exekveras som en SIMT (Single Instruction Multiple Thread) arkitektur vilket innebär att Kernel/s körs som ”lättviktstrådar” (många trådar körs åt gången och är lättvikts trådar (kan skapas och förstöras ”fort”) av en instruktion från CPU:n. Alla trådar kör samma Kernel kod, trådar har sitt eget ID för att separera dem från varandra. Trådar är grupperade i olika ”tråd-block”. Trådar kan kommunicera med andra trådar i samma tråd-block genom delat minne, trådar i samma tråd-block kan synkronisera exekvering. Trådar kan inte kommunicera med trådar i andra tråd-block på vår hårdvara. [39]

4.4.1 Synkrona och asynkrona CUDA exekveringar

CUDA exekveringar är asynkrona detta innebär att kontrollen återgår ”direkt” till CPU tråden som kallade på CUDA operationen. Detta är problematiskt om vi ska ta tiden med CPU tråden eftersom tiden som visas kommer inte vara tiden för själva CUDA operationen att exekvera utan endast tiden att ge GPU:n instruktioner. För att undvika detta följs teorin i Cuda C programming Guide [38] som säger att vi ska synkronisera CPU-tråden med CUDA exekveringen vid tidtagning. Detta gör vi under alla experiment även om vi inte påpekar det.

Listing 1: Synkronisering av CPU och GPU

1 cudaThreadSynchronize();

[38]

4.5 Accelererade bibliotek och THRUST

Det finns en mängd ”accelererade” bibliotek till CUDA C/C++ som möjliggör pro-grammerare att förbipassera lågnivåprogrammering i CUDA C/C++ inom områden som t.ex. matematik, strömningsmekanik, bild bearbetning m.fl. genom att använ-da dessa vältestade lösningar kan programmerare flytta fokus till resten av pro-gramvaran. En del av dessa accelererade bibliotek finns på Nvidia’s hemsida. Många är gratis men en del kostar (vissa ges ut i begränsade utgåvor utan betalning t.ex. CULA).

4.5.1 THRUST

Vi använder ett CUDA bibliotek som heter THRUST och försöker efterlikna C++ bib-lioteket ”Standard Template Library” (STL). THRUST är inkluderat i Nvidia CUDA 4.0 och nyare versioner. För att implementera THRUST funktioner i vår källkod används inkludera direktivet (med relevanta instruktioner) likadant som med STL. För att anropa THRUST funktioner skriver vi ”thrust” som namnrymd istället för t.ex. ”std”. [12]

Sorterings funktioner som THRUST använder har en komplexitet på O(nlogn) när ”mergesort” används och O(n) när ”radix” sortering används. THURST väljer radix

sortering vid inbyggda data-typer (float, int) och mergesort när radix sort inte kan användas. [28], [46].

Att implementera egna sorteringsalgoritmer är utanför omfattningen av denna stu-die men det finns en intressant artikel av Assarsson och Sintorn [3] som visar ett sätt att implementera sortering med CUDA och resultat av implementationen. Vi använder även THRUST till vektor inventering och binärsökning på GPU:n, det finns fullt stöd för dessa operationer [12].

4.5.2 CUBLAS

För matris-matris operationer använder vi här CUBLAS som står för ”Cuda Basic Linear Algebra Subprograms”. Problemet med att skriva egna GPU applikationer (Kernels) som utför matris-matris beräkningar är optimeringen av dem. Om vi skri-ver egna matris-matrismultiplikationer är det även redundant eftersom det är ett standard problem som är återkommande i datavetenskap. Även om vi inte behö-ver lika ingående kunskaper för CUBLAS som vi hade (För att utföra matris-matris applikationer) i CUDA C/C++ är det viktigt att förstå grunderna inom t.ex. GPU-struktur, minneshantering, trådar, Kernels (synkronisering etc) för att kunna få korrekta värden och dra rätt slutsatser.

Nvidia skriver att de i tester fått upp till 7 gånger snabbare prestanda än CPU/P med Intels MKL (Math Kernel Library) bibliotek som är ett ej kostnadsfritt, pro-prietär bibliotek[40] [41]. Prestandamätning på Eigens hemsida visar att den nya versionen (3.0.0) har liknande prestanda vid matris-matris multiplikation som In-tels MKL. [9]

5

Experiment

I denna studie utförs två sorters experiment. Det första experimentet handlar om vektorer och sortering, binärsökning, vektorinventering av dessa. Det andra expe-rimentet handlar om matrismultiplikation.

Först skrivs program som behandlar stora mängder individuella dataelement från ett exakt antal filer på hårddisken (8). Vi applicerar teorin praktiskt och samlar mätdata att analysera. Att behandla stora datamängder (bestående av många in-dividuella element) är ett problem som går att parallellisera och bemöta med det accelererade biblioteket THRUST och C++ biblioteket STL. Precis 8 filer behandlas vilket ger CPU:n bästa möjlighet att uppnå maximal parallellism eftersom STL algo-ritmerna är sekventiella exekverade på 4 parallella kärnor (8 emulerade trådar). Ur detta perspektiv uppnås hundra procent parallellism jämfört med de ursprungliga seriella programmen. Matrismultiplikation program är den andra delen av experi-menten. Slumpmässiga matriser genereras lokalt på processorns RAM minne och GPU:ns GDRAM minne. Matriser som behandlas är kvadratiska i stigande ordning (2, 16, 128, 256, 512, 2048, 4096).

Det är viktigt att mäta tiden på ett korrekt sätt som teorin påpekar. För att ta högupplöst tidmätning används windows ”QueryPerformanceCounter” [29] i samt-liga experiment.

t = (tklockaef ter− tklockaf ¨ore) (3)

Tiden för exekvering av en operation får vi med formeln ovan praktiskt implemen-terad. Variabeln tklockaf ¨ore anger vi värdet 0 vid tidmätningen.

För att minimera felkällor kontrollerar vi i aktivitetshanteraren att alla systemre-surser som kan påverkas allokeras åt experimenten. Med de parallella programmen används prestanda fliken i aktivitetshanteraren för att verifiera att samtliga 8 trå-dar och 4 kärnor allokeras till vårt exekverade program.

Figur 3: Bilden visar aktivitetshanteraren vid exekvering av den parallella delen av CPU vek-torprogrammet.

Figur 4: Bilden visar antalet trådar som används vid exekvering av CPU vektorprogrammet i underfönstret systemresurser. Average CPU har inte värdet 100.00 men detta är inte ett problem eftersom värdet gäller även tiden för applikationen att köra igång, läsa in filerna vi genererat samt skriva ut data till ett cmd fönster.

Figur 5: Bilden visar ströminställningen som används.

Bilden med ströminställningar visar att vi väljer hög prestanda inställningen vilket gör att CPU:n får spendera mer ström. CPU:n ökar frekvensen (speedUP) och valet gör vi får maximal tillgänglig prestanda. I samtliga experiment i denna studie används samma inställning.

5.1 Vektorgeneratorn

Vi skriver en generator som skapar olikt stora fält fyllda med slumpdata. Genera-torn skapar ett fördefinierat antal N fält av datatypen flyttal (float), i detta experi-ment 8 st. Fälten fylls med fördefinierat antal slumptal M med radbrytning mellan dem, inom denna studie 1024, 16384, en miljon och fem miljoner element. Ex-empel på dessa tal ur en fil är 5.411, 1199.485, 2153.940. På fälten utförs olika vektoroperationer med THRUST och STL.

Figur 6: Flödesschema på programmet som skriver ut stora fält fyllda med slumpmässiga flyttal till filer på hårddisken. Genom att definiera N som 8 får vi en fil/processor tråd.

5.2 Struktur på vektorhanterings program

CPU-seriella och CPU-parallella vektorhantering programmen beskrivs översiktligt av flödesschemat under.

Figur 7: Bilden visar översiktligt algoritmen som används till CPU vektor programmen.

Start är början på subrutinen som hanterar filerna som genereras. Filerna läses in till vars en individuell vektor. CPU operationer bearbetar (binärsöker, inventerar och sorterar) vektorerna. Vid varje vektor händelse tas tiden och lagras i en separatt vektor s.k. timeToABC på detta sätt har vektorn timeToABC tider vi är intresserade av. Vid varje iteration sker dessa operationer på varje fil som läses in till en vektor. För att veta tiden som något behöver att exekveras parallellt på CPU:n tas ur kloc-ka vektorn det största elementet. Alla andra mätningar måste ha avslutats innan därför ger största elementet totala tiden för hela vektoroperationen i parallella CPU subrutiner.

Listing 2: Vektor operationer

vector<float> timeToABC(NumberOfFiles);

for(int f=0; f<NumberOfFiles; f++){ take_time tn; . . . . vektoroperation pa CPU . timerTo = tn.get_time(); timeToABC.push_back(timerTo); }

Med GPU:n måste vi enligt teorin skicka data till och från grafikkortet manuellt. Vi måste även synkronisera GPU:n med CPU:n vid tidtagning på en GPU händelse med en CPU klocka, för att ta korrekt tid.

Figur 8: Bilden visar översiktligt algoritmen som används till GPU vektor programmen.

Listing 3: Skicka data till GPU

cudaThreadSynchronize(); take_time t0;

thrust::device_vector<float> gpuVector=cpuVector; cudaThreadSynchronize();

timerToSend= t0.get_time();

TimeToSendDataToDevice.push_back(timerToSend);

Vi tar inte hänsyn till hur CPU:n ska allokera resurser på GPU:n när THRUST används. Men vi måste säga till GPU:n vilken grafikkorts vektor vi menar d.v.s. måste vi sköta minneshantering men vi måste inte gå in på detaljer med trådar, trådblock etc. Vi allokerar och skickar data till en vektor på GPU:n med THRUST genom ett lika med tecken mellan GPU vektorn och CPU vektorn. Detta fungerar eftersom THRUST är programmerat att vara kompatibelt med STL.

Listing 4: Hämta data från GPU

cudaThreadSynchronize(); take_time t4;

thrust::copy(gpuVector.begin(), gpuVector.end(), cpuVector.begin()); cudaThreadSynchronize();

timerTo4 = t4.get_time();

TimeToReciveDataFromDevice.push_back(timerTo4);

För att hämta tillbaka vektorn till CPU:n (för att kunna läsa vektorn) måste vi ange tre parametrar till THRUST ”copy” funktionen. Dessa tre parametrar är början på vektorn vi vill skicka, slutet och början på CPU vektorn som ska ta emot innehållet.

5.3 Parallellt på CPU med OpenMP

De parallella implementationerna av CPU programmen är inkapslade av en ”nowait” OpenMP deklaration på detta sätt exekveras alla trådar samtidigt. I övrigt ser de parallella delarna av CPU programmen ut på samma sätt som de seriella därför behöver vi inte separat beskriva de olika vektor-operationerna för de parallella och seriella programmen.

Listing 5: Parallellt på CPU med OpenMP #define NumberOfFiles 8

. .

char startName =’a’;

omp_set_num_threads(NumberOfFiles);

#pragma omp parallel for shared(startName)

for(int j=0; j<NumberOfFiles; j++){ . .

. }

Koden ovan visar att NumberOfFiles definieras till 8 (vi har 8 emulerade kärnor). Denna tilldelning används som parameter till antalet trådar som ska användas vid OpenMP direktivet. For loopen inkapslas av ett OpenMP direktiv där argumentet startName används för att skilja de olika trådarna åt. Bokstäverna a och b är enda skillnaden i namnet på den första och andra filen vi läser in.

5.4 Sortering med THRUST och STL

För att sortera de slumpmässiga elementen används STL sorterings funktion och THRUST motsvarighet på GPU:n. Listorna sorteras till stigande ordning, vektorn som sorterings funktionen returnerar har tal från minsta till högsta ordning istället för den tidigare slumpmässiga ordningen.

Listing 6: Sortering på CPU STL’s sort

take_time t1;

std::sort(vecCPU.begin(), vecCPU.end()); timerTo = t1.get_time();

timeToSort.push_back(timerTo);

Koden ovan skickar argument till sorterings funktionen som början på en vektor (vecCPU) och slutet på en vektor. För att sortera på GPU:n kan vi använda THRUST på liknande sätt. Klocka vektorn heter timerTo.

Listing 7: Sortering på GPU med THRUST

cudaThreadSynchronize(); take_time t1; thrust::sort(gpuVector.begin(), gpuVector.end()); cudaThreadSynchronize(); timerTo = t1.get_time(); timeToSort.push_back(timerTo);

i koden ovan implementerar vi THRUST sortering på en osorterad vektor gpuVec. För att få korrekt resultat på tidtagningen är det viktigt att vi som teorin föreslår synkroniserar GPU med CPU vid tidtagning. Om vi inte gör detta kommer vi att ta tiden CPU:n behöver för att skicka instruktioner till GPU:n och sedan fortsätta ner i koden detta tar inte alls lika lång tid.

5.5 Binärsökning med THRUST och STL

Vi kontrollerar ifall ett element (324.4567) finns i sorterade vektorn cpuVec re-spektive gpuVec med binärsökning. Men binärsökningen sker på en redan sorterad vektor (sorterad i förra avsnittet).

Listing 8: Binär sökning på CPU med STL

take_time t2;

bool findVal = std::binary_search(vecCPU.begin(), vecCPU.end(), 324.4567); timerTo = t2.get_time();

timeToBinarySearch.push_back(timerTo);

Om flyttalet hittas i vektorn kommer variabeln findVal få värdet 1 annars 0. Argu-menten vi skickar in beskriver var sökningen ska börja, sluta och vilket tal vi letar efter.

Listing 9: Binär sökning på GPU med THRUST

cudaThreadSynchronize(); take_time t2;

bool findVal = thrust::binary_search(gpuVector.begin(), gpuVector.end() ,342.4567);

cudaThreadSynchronize(); timerTo2 = t2.get_time();

På GPU sidan anropar vi funktionen binary_search i thrust med samma typ argu-ment. Som vi ser är det relativt enkelt att implementera binärsökning med THRUST.

5.6 Vektorinventering med THRUST och STL

Vektor inventering med STL på CPU:n implementeras med följande kod.

Listing 10: Vektor inventering med STL på CPU

take_time t3;

std::reverse(vecCPU.begin(), vecCPU.end()); timerTo = t3.get_time();

timeToReverseVector.push_back(timerTo);

Nu är vektorn returnerad i omvänd ordning med hjälp av CPU:n och tid är tagen på detta. På GPU sidan kan vi till använda THRUST igen.

Listing 11: Vektorinventering på GPU med thrust

cudaThreadSynchronize(); take_time t3; thrust::reverse(gpuVector.begin(), gpuVector.end()); cudaThreadSynchronize(); timerTo3 = t3.get_time(); timeToVectorReverse.push_back(timerTo3);

5.7 Matrismultiplikation av kvadratiska matriser

Den andra delen av experimenten handlar om matrismultiplikation. Denna del är oberoende av vektorhanterings programmen. Vi genererar matriser i en stigande ordning, matriserna innehåller slumpmässiga flyttal. Två matriser A och B multi-pliceras till C biblioteken CUBLAS och Eigen används på GPU:n respektive CPU:n. 5.7.1 Matrismultiplikation på CPU

Matrismultiplikation på CPU utförs parallellt och seriellt.

Listing 12: MatrisMultiplikation på CPU:n med Eigen #include "Eigen\Dense"

using Eigen::MatrixXd; ...

const int size = 256; MatrixXd a(size,size); MatrixXd b(size,size); ...

//fyll a och b med slumpdata ...

MatrixXd res(size,size); omp_set_num_threads(8); res = a*b;

Vårt parallella CPU program har en OpenMP deklaration ”omp_set_num_threads(8)”. Den seriella versionen är identisk fast saknar OpenMP deklarationen. Matriserna vi multiplicerar är a och b, resultatet lagras i matrisen res. Vi tar tid på matrismul-tiplikationen genom att ta tiden innan och efter. Matris storleken ändras för att ta fram olika mätdata och varje ändring skiljs åt med ”Test ID”.

5.7.2 Matrismultiplikation på GPU med CUBLAS

Vi använder i denna studie biblioteket CUBLAS för matrismultiplikationer på GPU. CUBLAS ger oss kvalitativ kod på vår GPU. Precis som innan synkroniserar vi CPU med GPU vid tidtagning för att få korrekta mätdata.

Till matrismultiplikation använder vi CUDA C mer tydligt. Vi måste allokera minne på grafikkortet det gör vi nu med cudaMalloc. Vi måste kopiera den slumpmässiga datan till grafikkortet det gör vi med cudaMemcpy.

Listing 13: Allokera och skicka data till GPU:n

cudaMalloc((void**) &deviceA, memorySizeA); cudaMalloc((void**) &deviceB, memorySizeB); cudaMalloc((void**) &deviceC, memorySizeC);

cudaMemcpy(deviceA, hostA, memorySizeA, cudaMemcpyHostToDevice) ; cudaMemcpy(deviceB, hostB, memorySizeB, cudaMemcpyHostToDevice) ;

Hur mycket att allokera (memmorySize) räknar vi ut genom att multiplicera höjden och bredden för matriserna och sedan multiplicera resultatet med storleken av flyt-tal. Vi behöver inte skicka deviceC till grafikkortet eftersom den är tom.

Vi genererar två matriser separat (hostA, hostB) fyllda med slumpdata på CPU:n. Den slumpmässiga datan kopieras till grafikkortet (deviceA resp B) med cuda-Memcpy funktionen.

Listing 14: MatrisMultiplikation på GPU:n med CUBLAS

..

cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, deviceB, m, deviceA, k, &beta, deviceC, k);

Funktionen cublasSgem utför en matris-matris operation och matrisen C som ska-pas blir

C = αop(A)op(B) + βC (4)

CublasSgem är funktionen i CUBLAS som används för ”single-precision” matris-multiplikation på GPU:n och är förkortat för ”singleprecision general matrix mul-tiply”. Där single innebär att innehållet är i 32bitar (4byte) [49]. Vi måste ange en del argument till funktionen. Exakta API definitionen av dessa finns definerade i Nvidia dokumentet ”CUDAToolkit 4.1 CUBLAS Library” [40].

Kontexten handlar om b.la. vilken GPU som CUBLAS ska exekveras på och på vil-ken CPU-tråd i en multitråd applikation (samtliga CUBLAS funktioner behöver en handle och denna skapas med cublasCreate() funktionen).

Andra och tredje parametern säger att vi inte ska utföra transponent matris opera-tioner därför kommer op(A) och op(b) att kvarstå som A resp B. m anger antal rader matris A och C har. n anger antal kolumner matris B och C har. Det andra m anger antal kolumner Matris A har och antal rader matris B har. Alpha är skalären till ma-tris A i detta fall 1. deviceB pekar mot minnet vi allokerar för B. m är första indexet i matris B. deviceA pekar mot minnet vi allokerar för A. k efter deviceA är första index i matris A. Beta är en skalär som används till multiplikationen. Om Beta är noll är C inte ”valid input” (vilket den inte är i vårt fall). deviceC lagrar resultatet av matrismultiplikationen. Sista k är första index i deviceC matrisen. Nu är matris-multiplikationen färdig men vi måste utföra några steg till innan experimentet är färdigt.

Listing 15: Skicka tillbaka data till CPU

cudaMemcpy(host_res, deviceC, memorySizeC, cudaMemcpyDeviceToHost);

Nu kopieras resultatet av matrismultiplikationen från deviceC till host_res på CPU:n därför används ”deviceToHost” som den sista parametern i cudaMemcpy funktio-nen. Slutligen släpps allokerar minne fritt med free (CPU) funktionen och cudaFree (GPU) funktionen. Med denna kod utför vi matrismultiplikationen på GPU och an-tecknar mätdata.

6

Diskussion

Mätdata från experimenten analyseras här. Mätdata presenteras med hjälp av tabel-ler och stapeldiagram. Olika experiment skiljs åt av kolumnen ”Test ID”. Fullständig mätdata finns bifogade som appendix A resp B.

6.1 Analys av vektor program

Det är viktigt att vi är medvetna om att parallella CPU vektorprogrammen har fått en maximal parallellism (inget är seriellt) p.g.a. tekniska implementationen. Om algoritmen som utför t.ex. sortering parallelliserats hade detta inte nödvändigtvis varit fallet då en del av koden kanske inte går att utföra parallellt.

Snabbarekod=

1

0 +(1−0)₄ = 4 (5)

CPU vektorprogrammen bör bli ungefär 4 gånger snabbare där HT inte kan an-vändas. Om experiment visar att parallella CPU programmen är mer än 4 gånger snabbare beror det troligtvis på användning av HT teknologi enligt teorin. Hårdva-ran i denna studie har endast 4 riktiga kärnor.

Tabeller här innehåller mätdata från experimenten med vektorhanteringsprogram-men på CPU:n respektive GPU:n. Varje vektor-experivektorhanteringsprogram-ment involverar 8 filer seriellt på CPU:n, parallellt på CPU:n och på GPU:n. Vektorer sorteras med 1024, 16384, en miljon och 5 miljoner element. Genom att börja med små vektorer och sen öka i storlek observerar vi hur mycket operationer blir snabbare på GPU:n beroende på hur mycket data som skickas in. Vi letar efter en ungefärlig nedre gräns på antal data-element som behövs skickas in till GPU:n för att få snabbare programvara jämfört med CPU programmen.

Tabell 6, 7, 8 i appendix A innehåller mätdata för experimenten med CPU/S, CPU/P och GPU vektorprogrammen. Mätdata i tabell 8 gäller THRUST funktioner på GPU:n exkluderat tiden att skicka data till och från GPU:n. Tiden för minneshantering be-handlas i tabell 9. Genom att dividera värden från tabell 6, 7, 8 räknar vi ut antalet gånger GPU:n, CPU/P blev snabbare än CPU/S. Resultat av divisionen finns i ta-bell 3 detta resultat är inte räknat med tiden att skicka data fram och tillbaka till GPU:n.

Tabell 3: Resultat för olika vektorhanterings subrutiner exkluderat tiden att skicka data till och från grafikkortet, x betyder antal gånger snabbare.

CPU/P vs CPU/S GPU vs CPU/S GPU vs CPU/P Operation Storlek Test ID

6.4x 1.078x 0.168x Sortering 1024 1, 5, 9 4.925x 20.539x 4.17005x Sortering (16384) 2, 6, 10 4.168x 198.380x 47.589x Sortering 106 3, 7, 11 3.849x 293.690x 76.291x Sortering 5(106) 4, 8, 12 6.808x 1.366x 0.2006x Stabil sort 1024 1, 5, 9 5.653x 27.112x 4.795x Stabil sort (16384) 2, 6, 10 4.272x 225.151x 52.69x Stabil sort 106 3, 7, 11 4.128x 365.205x 88.456x Stabil sort 5(106) 4, 8, 12 0.08274x 0.0015x 0.01899x Vektor inv 1024 1, 5, 9 3.922x 13.357x 3.406x Vektor inv (16384) 2, 6, 10 4.111x 35.934x 8.741x Vektor inv 106 3, 7, 11 3.831x 179.808x 46.939x Vektor inv 5(106_{) 4, 8, 12} 4.233x 0.013x 0.0033x Binärsök 1024 1, 5, 9 5.680x 0.013x 0.0244x Binärsök (16384) 2, 8, 10 5x 0.013x 0.002x Binärsök 106 3, 7, 11 5.387x 0.014x 0.003x Binärsök 5(106_{) 4, 8, 12}

Genom att dividera tidsåtgången för tiden på de olika operationerna i tabellerna i appendix A (6 7 och 8) individuellt har vi fått fram ovanstående tabell. Kolumnen Test ID visar här vilka Test ID som divideras.

Amdhals lag förutspår väl hur mycket CPU/P programmen blir snabbare än CPU/S programmen när större mängder data behandlas. Men underskattar antalet gånger koden blir snabbare på de små datamängderna vid sortering och stabil sortering. Detta kan vara fallet eftersom HT kommer in i spel vid de små datamängderna. Vektor inventering verkar bete sig mer parallellt vid större vektorer och väldigt då-ligt vid små storlekar kan det bero på för små problem d.v.s. tar tiden för OpenMP biblioteket att allokera resurser längre än själva problemlösningen. Binärsökning verkar inte gå att behandla snabbt på GPU:n med THRUST. Men eftersom data tar relativt lång tid att skicka mellan GPU:n och CPU:n (se tabell 9) och tiden att bi-närsöka är kort är det bra att funktionen är tillgänglig med THRUST på GPU:n. Som vi ser i resultat tabellen (se tabell 3) har vi snabbare kod med THRUST på sorteringen parallellt först när vi använder vektorer som har ca 16 tusen element. Någonstans mellan 1024 och 16 tusen element finns det en gräns där GPU:n blir mer fördelaktig att använda än de parallella CPU programmen. Problemet är att även om GPU:n blir fördelaktig att utföra operationer på måste vi skicka data fram och tillbaka till GPU:n och den tiden är inte inkluderad i tabellen. Detaljerad tid för att skicka data till och från grafikkortet i vektoroperationerna finns i tabell 9 i appendix A.

stora datamängder över de seriella och parallella CPU programmen. Detta skulle sluta ske vid någon övre gränns p.g.a. hårdvarurestriktioner. Om vi försöker skicka in vektorer som är 1.5 GB stora skulle vi först vara tvungna att dela upp vektorn för att sen kunna skicka in den i mindre delar. Även om det inte skulle ta lång tid att dela vektorn skulle grafikkortet som används här aldrig kunna ta in mer än 1GB data åt gången p.g.a. hårdvarurestriktioner i dessa experiment. Det finns samma gränser med CPU:n även om dessa ibland är lättare att lösa genom att t.ex. uppgradera RAM minnet.

Parallella implementationerna av CPU programmen (se tabell 3) är ungefär 4 gånger snabbare än de seriella när mer data skickas in vilket är ungefär vad Amdhals lag förutspår. Att de ibland är mer än fyra gånger snabbare kan bero på att HT (Hyper Threading) kommer till nytta ibland.

Binärsökningen är enligt våra experiment snabbare på CPU:n än GPU:n (se ta-bell 3). Men om vi hadde inkluderat tiden att först sortera vektorerna på CPU:n så hade detta blivit annorlunda enligt mätdatan. Vi tar endast tiden att binärsöka och jämför.

Tabell 4: Total tid resultat. Tiden är summerad för de olika operationerna och därefter jämförd.

CPU/P vs CPU/S GPU vs CPU/S GPU vs CPU/P Element/fil Test ID

6.588x 0.716x 0,108x 1024 1, 5, 9

5.285x 11.074x 3.760x 16384 2, 6, 10

4.221x 38.909x 9.217x 106 3, 7, 11

3.997x 49.457x 12.373x 5(106) 4, 8, 12

Tiden summerad för de fyra vektor operationerna (minneshantering tiden inklu-derad för GPU:n) för varje Test ID ger oss tabell 10. Genom att dividera värden i total tid tabellen i appendix A fås tabellen ovan (tabell 4). Denna data kan vi presentera som ett stapeldiagram (se figur 9) för att få en översiktlig bild av prestandaökningen på GPU:n.

Figur 9: Diagramet visar GPU och CPU/P (CPU/S representerar ett). Diagramet är konstruerat efter total tid resultat tabellen.

GPU programmen är för en miljon element nästan 40 gånger snabbare och näs-tan 50 gånger snabbare för programmet som behandlar 5 miljoner element/vektor. Om vi inte ger GPU:n tillräckligt med data att bearbeta finns det ingen nytta av att använda GPU:n över CPU:n. Redan vid 16 tusen element ser vi att GPU:n presterar snabbare än det likvärdiga parallella CPU programmen.

Över parallella CPU programmen finns det först en nytta att använda GPU:n med 16384 element per vektor relativt CPU/P. Då blir GPU programmen 4 gånger snab-bare än CPU/P programmet. Någonstans mellan 1024 och 16384 data element per vektor går GPU:n förbi CPU/P programmet det blir nedre gränsen för totala tiden. GPU programmet är 9 gånger respektive 12 gånger snabbare än en miljon respek-tive fem miljoner parallella CPU/P programmet. Men det finns en övre gräns när detta hade slutat ske precis som för de individuella operationerna diskuterat ovan. Överlag är dessa resultat positiva för GPU:n om det finns tillräckligt med data att behandla parallellt. Där tillräckligt med data i detta fall verkar vara mer än 16000 element per vektor. Vid dessa GPGPU operationer är slutsatsen nådd av Langdon i performing with CUDA [26] att data är dyrt men beräkning tidsmässigt billig mat-chande med resultaten av våra experiment.

6.2 Analys av matrismultiplikation

experimen-data insamlad vid experimenten.

Tabell 5: Resultat Matrismultiplikation Eigen och CUBLAS (inkl tid att skicka ta emot data)

CPU/P vs CPU/S GPU vs CPU/S GPU vs CPU/P Test ID

0.294x 0.011x 0.038x 13 0.3x 0.032x 0.108x 14 0.445x 3.580x 8.041x 15 1.225x 23.264x 18.988x 16 4.118x 29.693x 7.20x 17 2.298x 27.692x 12.048x 18 2.013x 26.087x 12.95x 19 1.946x 29.379x 15.093x 20

Parallella CPU programmen blir upp till 4 gånger snabbare jämfört med de se-riella CPU programmen. Men vi kan se att CPU/P prestandan inte håller sig lika högt relativt CPU/S för matriser större/mindre än 512x512 och bearbetas med Ei-gen. Detta beror sannolikt på inre strukturen av Eigen, vilket inte ska tolkas som negativt eftersom det kan vara en ”trade off” för att få optimeringen av CPU:n till att börja med.

Vi måste ha en matris av ordningen 256 innan vi ser någon nytta av att använ-da CPU/n parallellt. Denna parallellism är inte samma som i vektorprogrammen eftersom denna inte är seriella algoritmer exekverade parallellt på processortrådar. Algoritmer i Eigen är sannolikt parallella och kanske även dynamisk (beroende på storleken) vilket skulle medföra att Amdhals lag inte kan tillämpas med F som 0.

[H]

Vi ser en del intressant fakta ur vår matrismultiplikation. Det krävs en matris någonstans mellan 16 ordningen och 128 ordningen stor innan det finns nytta med att använda GPU:n över de parallella CPU programmen. Detta blir nedre gränsen på denna hårdvaran mot CPU/P programmen. Anledningen är som teorin påpekar att det krävs en del tid att allokera och skicka/ta emot data från grafikkortet därför är det inte lönsamt att utföra matrismultiplikation på GPU:N med för små matriser. Men på större matriser (på vår hårdvara större än 256 ordnings matriser) kan det bli väldigt värdefullt.

Nvidia visar resultat att är CUBLAS upp till 5.5 gånger snabbare än Intels MKL vid ”SGEMM” operationer som är motsvarig till Eigen biblioteket i denna studie [43]. Vi får betydligt bättre prestanda i vissa fall här (15 gånger snabbare än CPU/P på TEST ID 20). Det beror sanolikt på en av följande faktorer. Eigen biblioteket pre-sterar sämre än MKL, Nvidia använder en betydligt snabbare processor men inte lika relativt snabbare grafikkort.

Det är intressant att veta vilket medelvärde som denna studie producerat i ma-trismultiplikationen. Medelvärdet för GPU:n får vi genom att summera kolumnen GPU vs CPU/p och dividera med antal rader i tabellen. Svaret blir ungefär 9.30x vilket överstiger Nvidias lägsta resultat. Även om CPU:n kan bli dubbelt så snabb under samtliga experiment med något påkostat bibliotek finns det fortfarande nytta av att använda GPU:n.

7

Slutsats

I denna studie har vi undersökt följande frågor som alla relaterar till hur vi kan utnyttja GPU för olika standardproblem inom datavetenskap. Den första frågan vi fokuserade på var hur ett modernt grafikkort vanligtvis fungerar och hur vi kan manipulera grafikkortet till att utföra egna operationer samt vilka alternativ det finns för att programmera grafikkort och vad vi bör tänka på här. Vi fann här att för en utvecklare finns det hjälpbibliotek att tillgå som kan underlätta processen av att programmera grafikkortet till att göra några vanligt förekommande rutiner inom datavetenskap. I våra experiment fokuserade vi speciellt på vanligt förkommande rutiner såsom sortering, binärsökning, vektorinventering och matrismultiplikation. Vi fann här att vi kunde urskilja en märkbar ökning i prestanda när vi exekverade ovannämnda rutiner på GPU:n i jämförelse med när dessa exekverades på CPU. Men under förutsättning att tillräckligt mycket data skickades till GPU:n.

Studien visar att det finns en tydlig nedre gräns på hur mycket data som gra-fikkortet behöver få att behandla innan det finns fördel av att använda grafikkor-tet jämfört med processorn. Vid små vektorer (ca 1024 element) visar resultat av studien att det finns liten nytta att använda grafikkortet (Nvidia GT550M) i expe-rimenten, jämfört med processorn (Intel i7 2670QM) på operationerna sortering, binärsökning, vektorinventering. Vid större vektorer finns det mer nytta att an-vända grafikkort. Vektorer med 5 miljoner element behandlas ungefär 12 gånger snabbare på grafikkortet än parallella motsvarigheter på processorn (inklusive ti-den att skicka data till och från grafikkortet). Matrismultiplikationer exekverar upp till 30 gånger snabbare på grafikkortet jämfört med seriella CPU motsvarigheten vid 512x512 stora matriser som innehåller flyttal.

Studien kan användas som en introduktion och vägledning för utvecklare i proces-sen att lära sig utnyttja GPU-kapaciten på datorer. Speciellt föreslås att utvecklare använder arkitekturen CUDA och bibliotek såsom THRUST och CUBLAS som ett stöd i denna process. Mätdatan kan användas som besluts underlag till steget mot att använda GPU:n. Vi har tagit fram ett antal program där vi ser att dessa med fördel kan exekveras på GPU vid mer datarika fall. Vi fann vidare att det finns stöd för en programmerare i denna process i form av arkitekturer (CUDA) samt bibliotek (THRUST och CUBLAS). Denna studie kan användas som ett stöd i processen för att implementera framtida program som med fördel kan exekveras på GPU.

7.1 Vidare studier

Det är möjligt att exekvera så kallade samtidiga Kernels. Dessa Kernels kan tän-kas på som trådar på processorn. Experimenten som behandlar vektorer här skulle vara intressanta att bemöta med en sådan konfiguration. Det skulle även vara in-tressant att se en implementerad automatisk ”hybrid” mjukvara som använder re-sultaten från denna studie för att bemöta t.ex. matriser under en viss ordning med processorn men byta automatiskt till grafikkortet när en viss ordning överstigs.

Referenser

[1] ALTERA. Internal Memory (RAM and ROM) User Guide. 11.1 January 2012 [2] Amdahl, Gene (1967). "Validity of the Single Processor Approach to Achieving

Large-Scale Computing Capabilities"(PDF). AFIPS Conference Proceedings (30): 483–485.

[3] Assarsson Ulf, Sintorn, Erik: Fast Parallel GPU-Sorting Using a Hybrid Algo-rithm. Journal of Parallel and Distributed Computing, 68 (10) pp. 1381-1388. [4] Björklund Maria, Paulsson Ulf.(2003) Seminarieboken att skriva, Presentera

och opponera.

[5] cplusplus.com. Sökord: Binary search.

Hämtad från <http://www.cplusplus.com/reference/algorithm/binary_search/> Hämtad den 3 juni 2012.

[6] cplusplus.com. Sökord: Sort.

Hämtad från <http://www.cplusplus.com/reference/algorithm/sort/> Häm-tad den 3 juni 2012.

[7] cplusplus.com. Sökord: Stable Sort.

Hämtad från <http://www.cplusplus.com/reference/algorithm/stable_sort/> Hämtad den 3 juni 2012.

[8] cplusplus.com. Sökord: Vector Reverse.

Hämtad från <http://www.cplusplus.com/reference/algorithm/reverse/> Hämtad den 16 oktober 2011.

[9] Eigen biblioteket.

Hämtad från <http://eigen.tuxfamily.org/>. Hämtad den 5 februari 2012 [10] Golub, Gene H.; Van Loan, Charles F. Matrix Computations (3rd ed.).

Baltimo-re: Johns Hopkins. (1996). [11] Gameranx.com.

Hämtat från <http://www.gameranx.com/gallery/evolution-of-graphics-1991-2011> Hämtad den 5 juli 2012.

[12] Jared Hoberock and Nathan Bell. Thrust: A Productivity-Oriented Library for CUDA. GPU Computing Gems. 2011-08-22.

[13] Jared Hoberock and Nathan Bell. Thrust a library of parallel algorithms Nvidia Thrust 1.5.1.

Hämtad från<http://code.google.com/p/thrust/> den 25 jan 2012.

[14] Intel Corporation Enhanced Intel SpeedStep Technology for the Intel Pentium M Processor. White Paper. March 2004.

[15] Intel coorporation. Sökord: HyperThreading.

Hämtad från <http://www.intel.com/support/processors/sb/> Hämtad den 16 april 2012.

[16] Intel coorporation. Sökord: i7 2670QM.

Hämtad från < http://ark.intel.com/products/53469> Hämtad den 25 april 2012.

[17] jcuda.org.

Hämtad från <http://www.jcuda.org/> Hämtad den 6 jun 2011.

[18] E.KANDROT, J.SANDERS. CUDA by example-an introduction to general-purpose GPU programming. 2011 pearson Education, Inc.

[19] khronos.org.

Hämtad från <http://gpgpu.org/developer/legacy-gpgpu-graphics-apis> Hämtad den 5 maj 2012.

[20] khronos.org alternativ till GPGPU. khronos.org.

Hämtad från <http://www.khronos.org/opencl/> Hämtad den 5 maj 2012. [21] D.KIRK, W.HWU. Programming Massively Parallel Processors A Hands-on

Ap-proach. 2010. Morgan Kaufman, [22] Komplett webbbutik. Sökord: Nvidia.

Hämtat från <http://www.komplett.se/k/search.aspx?q=Nvidia+>. Hämtad den 30 maj 2011.

[23] Komplett webbbutik. Sökord: Grafikkort.

Hämtat från <http://www.komplett.se/k/kl.aspx?bn=10488>. Hämtad de 2 jun 2011.

[24] Heiko Koziolek. The Role of Experimentation in Software Engineering. Seminar “Research Methods”, Summer Term 2005. 7th July 2005

[25] William B. Langdon . Debugging CUDA. GECCO Companion, CIGPU 2011 workshop, Simon Harding et al. Eds., p415–422.

[26] William B. Langdon . Performing with CUDA. GECCO Companion, CIGPU 2011 workshop, Simon Harding et al. Eds., p423–430.

[27] Brain Marshall. How Microprocessors Work.

Hämtad från <http://www.howstuffworks.com/microprocessor.htm> Hämtad den 20 jun 2011.

[28] D. Merrill, A. Grimshaw, Revisiting sorting for gpgpu stream architectures, In Proceedings of PaCT. 2010, 545-546.

[29] Windows Dev Center Microsoft .

[30] Vahid Mosavat. Kungliga Tekniska högskolan.

Hämtad från <http://www.csc.kth.se/utbildning/kth/kurser/DD1343/datae07/forel6.html> Hämtad den 25 jan 2012.

[31] MSDN OpenMP no wait.

Hämtad från: <http://msdn.microsoft.com/en-us/library/c8bys6hz(v=vs.80).aspx> den 20 jan 2012.

[32] Murphy Stein. CUDA Basics. New York University.

Hämtat från <http://www.cs.nyu.edu/manycores/cuda_many_cores.pdf> Hämtad den 6 dec 2011.

[33] Nvidia cooporation CUDA 4.0 application programmable interface. Hämtat från <developer.nvidia.com> Hämtad den 12 dec 2011.

[34] NVIDIA Corporation. Architecture Introduction & Overview 2009 NVIDA Corpo-ration.

[35] Nvidia Coporation.NVIDIA CUDATM_{.CUDA C BEST PRACTICES GUIDE 4.1.}

Ja-nuary 2012.

[36] Nvidia Coporation.Nvidia. NVIDIA CUDA GETTING STARTED GUIDE FOR MICROSOFT WINDOWS. April 2012.

[37] NVIDIA Corporation .GT550M specifications. 2011 NVIDA Corporation. Hämtat från <http://www.geforce.com/hardware/notebook-gpus/geforce-gt-550m/specifications> 16 jan 2012.

[38] NVIDIA Corporation.CUDA C Programming Guide. Version 4.0. 2011 NVIDA Corporation.

[39] Nvidia Coporation.NVIDIA CUDATM_{.NVIDIA CUDA C Programming Guide.}

Ver-sion 4.1. January 2012.

[40] NVIDIA Corporation. CUDA Toolkit 4.1 CUBLAS Library CUBLAS Library . 2012 Nvidia Corporation.

[41] Nvidia coporation CUBLAS documentation V4.1.

Hämtad från <http://developer.nvidia.com/cublas> den 12 dec 2011. [42] Nvidia Cuda Zone.

Hämtad från <http://developer.nvidia.com/category/zone/cuda-zone> 10 april 2011.

[43] Nvidia Cuda Zone CUBLAS.

Hämtad från <http://developer.nvidia.com/cuda/cublas> den 12 dec 2011. [44] OpenMP.

Hämtad från <http://openmp.org/wp/openmp-compilers/> Hämtad den 20 jan 2012.

[45] PCI SIG. PCI-SIG DELIVERS PCI EXPRESS 2.0 SPECIFICATION. BEAVERTON, Ore. – January 15, 2007.

[46] N. Satish, M. Harris, M. Garland, Designing efficient sorting algorithms for ma-nycore GPUs, in Proceedings 23rd IEEE Int’l Parallel & Distributed Processing Symposium, IEEE Computer Society, Washington, DC, 2009.

[47] SGI.com. Sökning: STL.

Hämtad från <http://www.sgi.com/tech/stl/> Hämtad den 18 jan 2012. [48] G.SPARR. Linjär Algebra. 1997-11-01 Studentlitteratur AB.

[49] John Stensby. Single-precision floating-point format. Version 1.0. EE448/528. [50] Ulrich Drepper Red Hat, Inc. What Every Programmer Should Know About

Memory November 21, 2007.

[51] Johannes Weidl. The Standard Template Library Tutorial. Friday, 26. April 1996.

[52] Wikipedia. Sökord: Matrixmultiplication.

Hämtad från <http://sv.wikipedia.org/wiki/Matris#Matrismultiplikation> den 8 maj 2012.

[53] Wynters Erik. Parallel processing on NVIDIA graphics processing units using CUDA. Journal of Computing Sciences in Colleges archive Volume 26 Issue 3, January 2011. Pages 58-66.

A

Mätdata vektorhanterings program

Tid (ms) Operation Element/fil Test ID

8 std::sort 1024 1 10.048 std::stable_sort 1024 1 t<0.001 std:: reverse 1024 1 0.051 std::binary_search 1024 1 213 std::sort 16384 2 281.16 std::stable_sort 16384 2 8 std:: reverse 16384 2 0.583 std::binary_search 16384 2 16045 std::sort 106 3 18154 std::stable_sort 106 3 525 std:: reverse 106 3 0.08 std::binary_search 106 3 88533 std::sort 5(106_{) 4} 109905 std::stable_sort 5(106) 4 2627 std::reverse 5(106_{) 4} 0.09 std::binary_search 5(106_{) 4}

Tabell 7: STL operationer på CPU Parallell

Tid (ms) Operation Element/fil Test ID

1.251 std::sort 1024 5 1.476 std::stable_sort 1024 5 0.011 std:: reverse 1024 5 0.012 std::binary_search 1024 5 43.245 std::sort 16384 6 49.731 std::stable_sort 16384 6 2.040 std:: reverse 16384 6 0.103 std::binary_search 16384 6 3849.01 std::sort 106 7 4249.51 std::stable_sort 106 ₇ 127.71 std::reverse 106 7 0.016 std::binary_search 106 7 22998 std::sort 5(106_{) 8} 26620 std::stable_sort 5(106_{) 8} 685.78 std::reverse 5(106_{) 8} 0.017 std::binary_search 5(106_{) 8}

Tabell 8: Tiden att exekvera operationer på GPU:n med THRUST.

Tid (ms) Operation Element/fil Test ID

7.429 thrust::sort 1024 9 7.354 thrust::stable_sort 1024 9 0.572 thrust::reverse 1024 9 3.638 thrust::binary_search 1024 9 10.370 thrust::sort 16384 10 10.333 thrust::stable_sort 16384 10 0.599 thrust::reverse 16384 10 4.173 thrust::binary_search 16384 10 80.88 thrust::sort 106 ₁₁ 80.63 thrust::stable_sort 106 11 14.61 thrust::reverse 106 11 6.02 thrust::binary_search 106 ₁₁ 301.45 thrust::sort 5(106_{) 12} 300.94 thrust::stable_sort 5(106_{) 12} 14.61 thrust::reverse 5(106_{) 12} 6.2 thrust::binary_search 5(106_{) 12}

Tabell 9: Tid för minneshantering på GPU med THRUST

Tid att skicka data (ms) Tid att ta emot data (ms) Test ID

4.371 1.905 9

12.694 7.181 10

335.74 374.56 11

1616.06 1862.22 12

Tabell 10: Total tid

Total tid (ms) Enhet/Läge Test ID

18.099 CPU/S 1 502.218 CPU/S 2 34724.08 CPU/S 3 201065.09 CPU/S 4 2.747 CPU/P 5 95.027 CPU/P 6 8226.246 CPU/P 7 50303.797 CPU/P 8 25.272 GPU 9 45.351 GPU 10 892.44 GPU 11 4065.48 GPU 12

B

Mätdata matrismultiplikation program

CPU/S (ms) CPU/P (ms) Dimensioner Test ID

0.00233 0.00793 2x2 13 0.00699 0.02335 16x16 14 1.58985 3.5706 128x128 15 16.913 13.805 256x256 16 111.489 27.0698 512x512 17 553.221 240.683 1024x1024 18 3871.61 1923.19 2048x2048 19 31006.6 15930.2 4096x4096 20

Tabell 12: Matrismultiplikation på GPU

GPU (ms) Minnes hantering (ms) Dimensioner Test ID

0.0849 0.119 2x2 13 0.0970 0.119 16x16 14 0.222 0.222 128x128 15 0.388 0.339 256x256 16 2.1967 1.558 512x512 17 15.891 4.086 1024x1024 18 124.787 23.619 2048x2048 19 993.3 62.1015 4096x4096 20