Führen Sie den Take a look at einige Male aus, wir erhalten die gleiche Ausführungszeit über viele Iterationen hinweg, was bedeutet, dass etwas nicht stimmt oder wir zwischengespeicherte Antworten erhalten.
Und das ist die Timer-Funktion, mal sehen, ob sie funktioniert, sie tut es:
2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(1000) 2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1) Milliseconds (15)
Okay, wonach suchen wir?
Was auch immer parallel getan wird, wird in der Nähe der Zeit sein, oder? Wir suchen additionally nach dieser zusätzlichen Verzögerung, die im Vergleich zur Gesamtausführung nicht winzig ist
Beginnen wir additionally mit 2048 Elementen und erhöhen die Iterationen, damit wir mal bemerken können, dass wir es noch einmal mit 1 Million versuchen.
Lassen Sie uns tatsächlich 1 Artikel und 1 Million Iterationen senden.
2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(1000000) 2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1) Milliseconds (16)
Fantastisch, 1 Artikel 100 Millionen Iterationen
2023.05.03 02:51:17.223 blog_,benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(100000000) 2023.05.03 02:51:17.223 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 02:51:17.223 blog_benchmark_2 (USDJPY,H1) Milliseconds (16)
Okay, geänderter Kernel wieder zu diesem:
string kernel="__kernel void memtests(__global int* group_id," "int iterations){" "double sum=(double)0.0;" "double inc=(double)-2.6;" "double step=(double)0.01;" "double radi=(double)0.0;" "double rstep=(double)0.00001;" "for(int i=0;i<iterations;i++){" "sum=((double)tanh(inc))*((double)sin(radi));" "inc+=step;" "radi+=rstep;" "if(inc>2.6&&step>0.0){step=(double)-0.01;}" "if(inc<-2.6&&step<0.0){step=(double)0.01;}" "if(radi>6.28&&rstep>0.0){step=(double)-0.00001;}" "if(radi<0.0&&rstep<0.0){step=(double)0.00001;}" "}" "group_id(get_global_id(0))=get_group_id(0);}";
Auf 10000 Iterationen gesenkt, übertreiben Sie es im Allgemeinen nicht, da Sie möglicherweise Cache-Antworten erhalten und in dem Second, in dem Sie die Berechnung ändern, sagen wir, Sie werden Ihre GPU beschweren hören.
Additionally 10000 Iterationen:
2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(10000) 2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1) Milliseconds (31)
lass uns auf 100K erhöhen
Ich werde wieder 31 ms bekommen, da bin ich mir sicher, aber dieses Mal werde ich es auch nach einem Neustart testen.
Additionally 1 Merchandise 100K Iterationen kein Neustart des PCs.
2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(100000) 2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1) Milliseconds (47)
Okay und 1 Punkt 100K Iterationen Neustart des PCs.
2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(100000) 2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1) Milliseconds (47)
Gleichzeitig yay okay, wir sind intestine mit diesem Kernel FILALLY und schauen Sie, wer sich entschieden hat, aufzutauchen:
2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1) Kernel non-public mem (40)
die Messung des privaten Gedächtnisses hörte erstmals auf, Null zu sein. Okay, das sind additionally 40 was? 40 Bytes professional Artikel ?
Lassen Sie uns den Kernel messen, wir zählen nicht globale und nicht lokale und nicht Argumente, additionally:
ja 5 verdoppelt 5 mal 8 seine 40 Bytes. Okay additionally das geht auch. Eindrucksvoll.
Der non-public Speicher ist superschnell, Sie müssen sich keine Sorgen machen. Wie wir die Grenze messen, ist eine spätere Frage.
In Ordnung, kann es jetzt 1 Million Iterationen dauern?
2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1) Work Gadgets (1) Iterations(1000000) 2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1) Work Teams (1) 2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1) Milliseconds (203)
Okay, aber ist diese Zahl etwas, das es uns ermöglicht, die « Falte » zu erkennen?
Mal sehen, fangen wir an, die Artikel zu erhöhen, ich werde die Protokolle nicht einfügen, ich werde sie nur auflisten
| Artikel# | Gruppen# | MS |
|---|---|---|
| 1 | 1 | 203 |
| 100 | 1 | 203 |
| 256 | 1 | 188 |
| 512 | 2 | 297 |
| 1024 | 4 | 578 |
| 2048 | 8 | 1125 |
| 4096 | 1 | 2235 |
Hm, was ist da unten gelaufen? Wir gingen von 8 Gruppen zu 1 Gruppe und der Bildschirm blitzte auf
Additionally sind 1024 die maximalen Elemente professional Recheneinheit und 256 die maximale Gruppengröße, die 4 1024 gleichzeitige Prozesse und 4 Gruppen angibt, die im selben Speicher arbeiten können? (für 1D unter der Annahme, dass wir selbst keine Aufteilung vornehmen)
Wir müssen uns daran erinnern, dass wir es bei diesem Benchmark über die Splits entscheiden lassen und es hat 1 group ausgewählt, oder es hat hier nichts parallel getan und die von uns ausgegebene Datei bestätigt dies
GROUP.ID(4095)=0
der letzte Eintrag struggle 0 . Es hat additionally nichts berechnet und wahrscheinlich gelitten.
Sooo, wenn ich mir die obigen Zahlen anschaue, glaube ich nicht, dass es eine Parallelität gibt, wenn 4 Gruppen parallel ausgeführt würden oder 2 Gruppen oder 8 Gruppen, würden wir mehr oder weniger die gleiche Zeit mit 1 Gruppe sehen, oder? Was additionally parallel läuft, sind vielleicht die 256 Gadgets.
Hmm, additionally führen wir jetzt einen weiteren Parameter in unseren Take a look at ein und nennen ihn Benchmark 3 . Der lokale Dimensionsparameter im Ausführungsaufruf
bool CLExecute( int kernel, uint work_dim, const uint& global_work_offset(), const uint& global_work_size(), const uint& local_work_size() );
Was ist, wenn ich dies auf 32 setze? Wir können 8 Untergruppen erhalten.
Diese 8 Untergruppen werden gleichzeitig ausgeführt, aber wenn wir 10 Untergruppen (dh 2 Gruppen) bereitstellen, erhalten wir eine höhere Ausführungszeit?
Ich denke, ich muss das nicht testen, aber lassen Sie uns bestätigen.
Soweit ich weiß, habe ich ein Gerät, das 256 Elemente gleichzeitig berechnen kann, und deshalb die Ganzzahl
CL_MAX_WORK_GROUP_SIZE gibt diese Zahl auch zurück, die Gruppengröße kann 256 sein, weil 256, die gleichzeitig laufen, den lokalen Speicher teilen können.
Könnte es so einfach sein?
Lassen Sie uns testen, wie sich die Gruppen aufteilen, wenn ich den Parameter local_work_size auf 32 setze (passend zur Warp/Wellenfront) für verschiedene Anzahl von Elementen
uint work_items_per_group()={32}; lengthy msStarted=GetTickCount(); CLExecute(ker,1,offsets,works,work_items_per_group);
Ich glaube nicht, dass ich sonst noch etwas ändern muss.
Mal sehen, wie viele Gruppen es erstellt
(nebenbei bemerkt, der 4096 ist tatsächlich abgestürzt, da ich beim nächsten Lauf einen Kontextfehler bekam, der einen Neustart erforderte)
| Artikel | Gruppen |
|---|---|
| 1 | keine, Fehler natürlich okay |
| 32 | 1 (203ms) |
| 64 | 2 (203 ms) |
| 128 | 4 (203ms) |
| 256 | 8 (203 ms) |
| 512 | 16 (297 ms) |
| 1024 | 32 (578 ms) |
| 2048 | 64 (1125 ms) |
| 4096 | 1 (2234 ms) |
Okay, offensichtlich die
lengthy CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE);
die in meinem Gerät 1024 zurückgibt, bezieht sich auf die maximale Anzahl paralleler Elemente, die gleichzeitig ausgeführt werden können.
Das heißt, wenn Sie die lokale Größe auf 1024 setzen, werden 1024 Elemente nicht parallel ausgeführt, sondern in Pakete von 1024 aufgeteilt, und es wird wahrscheinlich fehlschlagen und langsamer werden, aber wenn Sie es auf 1 setzen, behandelt es jedes Arbeitselement als seine eigene Gruppe , Rechts ?
Wenn ich additionally für diesen Take a look at, der keine Gruppen benötigt, die lokale Größe auf 1 setze, bekomme ich die schnellste Ausführung, wenn die Elemente auf 1024 tot sind?
Nein, ich bekomme wieder einen Fehler. Was ist, wenn die Gesamtzahl der Artikel 256 beträgt?
nein, Glitch … ähm Vielfache von 2 vielleicht?
Nnö, Panne … additionally was um alles in der Welt ist da nochmal los?
Was ist, wenn ich 8 Artikel und 1 lokalen habe, additionally 1 Artikel professional Gruppe 8 Gruppen im Wesentlichen
Kein Fehler, additionally ist die Anzahl der Gruppen begrenzt, aber das hat nichts damit zu tun:
CL_DEVICE_MAX_WORK_GROUP_SIZE
Und die Möglichkeit, die Gruppengrenze zu ermitteln, besteht darin, nach dem Abfall von der inkrementierenden Anzahl von Gruppen auf 1 zu suchen, gefolgt von der Bildschirmstörung ???
Aber es hat auch 2048 Arbeitselemente toleriert, additionally sehen wir, glaube ich, die Nichtkommunikation zwischen OpenCL und der {Hardware}.
Ist es 8 professional Dimension, additionally 8^3? 512 ? oder die Kubikwurzel von 1024 (die maximalen Elemente ? ) hmmm
das sind ~10, additionally maximal 10 Gruppen? (professional Dimension nehme ich an)
Woher sollen wir das wissen? 320 Gadgets mit 32 Native würden wahrscheinlich funktionieren
Es struggle okay, aber 64 funktionierten zu 2048 Artikel mit 32 lokalen … sowieso … ich bin abgerissen, ich gap es ein anderes Mal ab.
Benchmarks hänge ich an…
Auch aus all diesen Checks kann man etwas mitnehmen, und zwar den Wert
CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE);
Sagt Ihnen, wie viele Kernel-Instanzen parallel laufen können Dieses Gerät (oder professional Recheneinheit, aber das kann mit der CPU getestet werden)