Performance

mw1039
Computerversteher
Computerversteher
Beiträge: 346
Registriert: 12. Apr 2011 12:18

Re: Performance

Beitrag von mw1039 » 1. Feb 2012 10:50

Lucas- hat geschrieben:cpu variante [...] 1250 ms
ThorbenKo hat geschrieben:OCLGPU 2392 millisec.[...]mir würde das so reichen
Die Zeiten sehen danach aus, als ob es noch nicht reichen wuerde.
Wenn man einen Thread pro Workgroup anschmeisst, missbraucht man die GPU hart. Die GPU ist kein! Multicoresystem, in dem man alle Threads voneinander unabhaengig betrachten kann. Man muss die SIMD-Eigenschaft der einzelnen CUDA-Cores ausnutzen, ansonsten nutzt man nur ein Zweiunddreissigstel der Rechenpower, die man zur Verfuegung hat.

ThorbenKo
Windoof-User
Windoof-User
Beiträge: 25
Registriert: 7. Nov 2011 19:16

Re: Performance

Beitrag von ThorbenKo » 1. Feb 2012 13:08

mw1039 hat geschrieben:
Lucas- hat geschrieben:cpu variante [...] 1250 ms
ThorbenKo hat geschrieben:OCLGPU 2392 millisec.[...]mir würde das so reichen
D.
Die 2392 ms waren für loop 10. Mit loop 10 braucht die CPU 10 * 1250 ms = 12500 MS und damit mehr als meine 2392ms, oder lieg ich da falsch?

JulM
Mausschubser
Mausschubser
Beiträge: 47
Registriert: 1. Okt 2009 19:25

Re: Performance

Beitrag von JulM » 4. Feb 2012 04:42

Wie wird denn genau der Speed-Up berechnet? Beim letzten Praktikum wurden hier ja so schöne Formeln gepostet, damit war der Speed-Up ja einfach auszurechnen.
Ich bin jetzt einfach davon ausgegangen das (execution time cpu) / (execution time gpu) = speed up, ist das richtig so?

Denn noch bekomme ich meinen Kernel nicht unter 126 ms :evil: , dabei habe ich mich sehr an den reinen CPU-Code gehalten. Um es noch etwas zu reduzieren müsste man wahrscheinlich, die Berechnung noch verändern oder?

Darf man denn fragen, welche Werte für die globalWorkSize und localWorkSize bisher die besten Ergebnisse gebracht haben, oder ist damit schon zu viel der Lösung verraten?

bttf
Mausschubser
Mausschubser
Beiträge: 76
Registriert: 9. Nov 2011 14:04

Re: Performance

Beitrag von bttf » 4. Feb 2012 14:48

(execution time cpu) / (execution time gpu) = speed up
Davon gehe ich auch mal aus :roll:

Wir haben im Grund auch nur den CPU Code parallelisiert. Der Speedup liegt mit 70 ms allerdings bei etwa 15,5 fach. Hast du auch nur die untere Dreiecksmatrix berechnet? :wink:

Wir haben nur die localWorkSize angepasst auf 7*7 Work Items. Damit wird dann meiner Meinung nach der Kernel-Code etwas aufgeräumter, weil man z.B. zur Berechnung der twoPointOccs dann nur eine Zeile braucht, weil 42 Work Items zur Verfügung stehen (die restlichen 7 Work Items von den 49 machen bei uns einfach nichts. Es ging uns nur darum 42 Work Items in einer Work Group zu haben, deshalb die localWorkSize von 7*7. Ob das das Ganze auch schneller macht....keine Ahnung :mrgreen: ).

mw1039
Computerversteher
Computerversteher
Beiträge: 346
Registriert: 12. Apr 2011 12:18

Re: Performance

Beitrag von mw1039 » 5. Feb 2012 10:38

JulM hat geschrieben:Wie wird denn genau der Speed-Up berechnet? [...] (execution time cpu) / (execution time gpu) = speed up, ist das richtig so?
Genau, solange ihr nur die reine Berechnungszeit ohne Dateiladen etc. beruecktsichtigt.
JulM hat geschrieben:Denn noch bekomme ich meinen Kernel nicht unter 126 ms , dabei habe ich mich sehr an den reinen CPU-Code gehalten. Um es noch etwas zu reduzieren müsste man wahrscheinlich, die Berechnung noch verändern oder?
Wahrscheinlich musst du dann wirklich noch etwas mehr vom CPU-Code abweichen. Es gibt auch oft Anwendungsfaelle, bei denen man auf der GPU ganz andere Wege geht als auf der CPU, damit die GPU voll ausgelastet wird.
Du kannst ja mal versuchen eine Art Profiling zu machen, indem du nacheinander einzelne Codeabschnitte (von einer Synchronisation zur naechsten) doppelt ausfuehren laesst (einfach durch Codekopieren). Dann ist zwar u.U. das Endergebnis nicht mehr korrekt aber du kannst sehen, welcher Abschnitt die Ausfuehrungszeit am meisten verlaengert.
bttf hat geschrieben:Es ging uns nur darum 42 Work Items in einer Work Group zu haben, deshalb die localWorkSize von 7*7. Ob das das Ganze auch schneller macht....keine Ahnung ).
Ihr koennt die localWorksize beliebig machen, z.B. auch 42*1. Das haette auch den Vorteil, dass ihr nurnoch auf get_local_id(0) zugreifen muesst.
Generell ist es sinnvoll, wenn die localWorksize ein Vielfaches der Zahl der Threads ist, die ein CUDA-Core gleichzeitig ausfuehrt, also 32. Wenn die localWorksize ein echtes Vielfaches (also nicht 1*32 sondern 2*32, ...) ist, kann der Cudacore auch Latency Hiding machen, d.h. waehrend er fuer einen Satz von 32 Threads auf Speicheranfragen wartet, kann er in der Zwischenzeit an einem anderen Satz von 32 Threads weiterrechnen.

philipp_m
Mausschubser
Mausschubser
Beiträge: 99
Registriert: 4. Dez 2010 18:10

Re: Performance

Beitrag von philipp_m » 5. Feb 2012 11:02

Bin jetzt auch bei einem Speedup von 16,3 gelandet, wichtig war echt der Tipp, nur die untere Dreiecksmatrix zu berechnen, da man das normalerweise bei der OpenCL-Variante nicht unbedingt direkt tut. Ich wüsste auf Anhieb noch einige Stellen, die sich optimieren lassen, aber so auf die Schnelle bin ich zufrieden:)

Ach übrigens, der Rechner 07 ist immer noch extrem langsam - ich glaube, auf diesem ist ein 10facher SpeedUp gar nicht möglich.

Noch ein paar Vergleichswerte, die eine relativ einfache Lösung auf Rechner 5 ausspuckt:
LOOPCOUNT = 100

Code: Alles auswählen

inpath: test.fasta
MI outpath: test.fasta.MI.out
computing on the CPU

number of sequences: 42
sequence length: 1010

computing MI
execution time: 110570 milliseconds
writing output to test.fasta.MI.out

Code: Alles auswählen

inpath: test.fasta
MI outpath: test.fasta.MI.out
computing on the GPU

number of sequences: 42
sequence length: 1010

computing MI
kernel created, ready for computation
execution time: 6780 milliseconds
writing output to test.fasta.MI.out

Code: Alles auswählen

 comparing pr5/cpu.out and pr5/test.fasta.MI.out
comparison of 510555 entries finished
maximum relative deviation was 0.00099813349038 %, average relative deviation was 7.59192111881e-06 %

neoatlant
Mausschubser
Mausschubser
Beiträge: 58
Registriert: 14. Okt 2010 22:16

Re: Performance

Beitrag von neoatlant » 5. Feb 2012 11:54

Wir haben nur die localWorkSize angepasst auf 7*7 Work Items. Damit wird dann meiner Meinung nach der Kernel-Code etwas aufgeräumter, weil man z.B. zur Berechnung der twoPointOccs dann nur eine Zeile braucht, weil 42 Work Items zur Verfügung stehen (die restlichen 7 Work Items von den 49 machen bei uns einfach nichts. Es ging uns nur darum 42 Work Items in einer Work Group zu haben, deshalb die localWorkSize von 7*7. Ob das das Ganze auch schneller macht....keine Ahnung :mrgreen: ).
verstehe ich es richtig wenn ich vermute dass du einen Workgroup zur Berchnung eines MI Eintrags benutzt ?

bttf
Mausschubser
Mausschubser
Beiträge: 76
Registriert: 9. Nov 2011 14:04

Re: Performance

Beitrag von bttf » 5. Feb 2012 12:26

Ja richtig. Eine Work Group berechnet bei uns immer genau einen Matrixeintrag.

Wie Michael Wächter oben geschrieben hat, hätte man auch eine localWorkSize von 42*1 nehmen können. Aber da wir als erstes im Code die Work Items "durchnummerieren" (int number_work_item = (get_local_id(0) * 7) + get_local_id(1);), um dann damit einfacher weiterzurechenen, kommt es etwa auf's gleiche raus :-)

Unser Ergebnis hat jetzt auch einen Speedup von 16,3... ich konnte noch mal unglaubliche 3 ms rausholen :mrgreen:

daehn
Erstie
Erstie
Beiträge: 16
Registriert: 5. Okt 2010 13:42

Re: Performance

Beitrag von daehn » 9. Feb 2012 16:33

Wir hatten bei 100 Ausführungen auf der GPU (Cluster 06) 1685ms insgesamt, die CPU hat im Vergleich bei uns 110930ms gebraucht. Also etwa 65 fach Speedup.
Wir haben Submatrizen zu je 4x4 Einträgen berechnet, 22 Threads pro einem solchen Eintrag.
Allerdings haben wir angenommen, dass NUMSEQ < 256 (statt 512).

mw1039
Computerversteher
Computerversteher
Beiträge: 346
Registriert: 12. Apr 2011 12:18

Re: Performance

Beitrag von mw1039 » 9. Feb 2012 18:02

Wow, gut.

Eine Sache, die man z.B. noch machen koennte, ist, dass man vom Kernel unterschiedliche Versionen hat (in CUDA mit Templates, in OpenCL koennte man das mit unterschiedlichen Versionen der Kernelcode-Datei realisieren) und abhaengig von der Sequenzanzahl eine unterschiedliche Codeversion benutzt. Mit wenigen Sequenzen nimmt man kleine Datentypen, mit mehr Sequenzen groessere und muss dann Geschwindigkeitsverlust in Kauf nehmen.

Antworten

Zurück zu „Archiv“