Performance

[mw1039]

cpu variante [...] 1250 ms

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]

cpu variante [...] 1250 ms

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]

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 , 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]

(execution time cpu) / (execution time gpu) = speed up

Davon gehe ich auch mal aus

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?

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 ).

[mw1039]

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.

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.

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]

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]

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 ).

verstehe ich es richtig wenn ich vermute dass du einen Workgroup zur Berchnung eines MI Eintrags benutzt ?

[bttf]

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

[daehn]

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]

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.