Hallo.
Kurzfassung:
1.CLGetInfoInteger(Kernel,CL_KERNEL_WORK_GROUP_SIZE); gibt Ihnen die Anzahl der Kernel-Instanzen, die gleichzeitig auf dem Gerät (oder der Recheneinheit) ausgeführt werden können (nur mit 1CU-GPU getestet). Es wird auch die maximal mögliche Anzahl von Elementen in einer Gruppe sein.
2.Wenn 1 Arbeitsmaß können Sie angeben, wie viele Arbeitselemente Ihre Arbeitsgruppen haben sollen In mit dem non-compulsory hervorgehobenen Array
bool CLExecute( int kernel, uint work_dim, const uint& global_work_offset(), const uint& global_work_size(), const uint& local_work_size() );
—————————–Lange Model
Dieser Take a look at stellt einen eindimensionalen « Arbeitsbereich » mit den verfügbaren nativen mql5-OpenCL-Befehlen bereit.
Die Aufgabe besteht darin , herauszufinden , wie lokaler Speicher Arbeitsgruppen zugeordnet werden kann .
Lassen Sie uns zunächst einen einfachen Take a look at erstellen und versuchen, abzuschätzen, was die GPU tut oder wie die GPU die Arbeitslast ohne Anweisungen aufteilt, wenn Sie so wollen.
Ich weiß nicht, ob es etwas bedeutet, aber das versuche ich herauszufinden.
Es muss auch immer und bei jeder Erkenntnis im Hinterkopf behalten werden, dass dies eine spezifische Reaktion dieser {Hardware} sein kann, sodass die Flexibilität vorhanden sein muss, damit andere (Sie) es selbst testen können.
OpenCL hat einige eingebaute Indizes, die vom Kernel aufgerufen werden können und die bereitstellen
- Globale ID des Arbeitselements für eine Dimension
- Lokale ID des Arbeitselements für eine Dimension
- Gruppen-ID des Arbeitselements für eine Dimension
Wir können additionally diese Indizes nutzen und sehen, wie die Gruppen angeordnet sind, wenn wir eine Reihe von Aufgaben auf der GPU ablegen.
dies ist der Kernel, die Indexaufrufe haben die in Klammern angegebene Dimension
string kernel="__kernel void memtests(__global int* global_id," "__global int* local_id," "__global int* group_id){" "global_id(get_global_id(0))=get_global_id(0);" "local_id(get_global_id(0))=get_local_id(0);" "group_id(get_global_id(0))=get_group_id(0);}";
Was dies tut, ist, dass es 3 Integer-Arrays des globalen Speichers global_id, local_id, group_id empfängt und sie mit dem relevanten Index an der globalen Place füllt. Wenn wir zum Beispiel 10 Orangen in 2 Tüten hätten, würden wir den Tütenindex dem linearen Array-Index von Orangen zuweisen.
Wir würden sagen, Orange(0) ist in Bag0 und Orange(9) ist in Bag1 , wir würden den Index der Orange in der Tasche nicht verwenden (Orange(0) ist in Bag0 und Orange(4) ist in Bag1). würde uns nichts darüber sagen, wie die Orangen angeordnet waren!
Daher die Verwendung von get_global_id(0) zum Speichern aller IDs.
Wir machen all die langweiligen und lästigen OpenCl-Vorbereitungskrams, wie das Programm läuft, wann es aufgerufen wird, puffern Kernel-Kontexte und wir stellen einen Parameter dafür bereit, wie viele Elemente oder Orangen zum Take a look at geworfen werden.
so sieht der Code aus:
#property model "1.00" int OnInit() { EventSetMillisecondTimer(33); return(INIT_SUCCEEDED); } void OnDeinit(const int cause) { } void OnTimer(){ EventKillTimer(); int ctx=CLContextCreate(CL_USE_GPU_DOUBLE_ONLY); if(ctx!=INVALID_HANDLE){ string kernel="__kernel void memtests(__global int* global_id," "__global int* local_id," "__global int* group_id){" "global_id(get_global_id(0))=get_global_id(0);" "local_id(get_global_id(0))=get_local_id(0);" "group_id(get_global_id(0))=get_group_id(0);}"; string errors=""; int prg=CLProgramCreate(ctx,kernel,errors); if(prg!=INVALID_HANDLE){ ResetLastError(); int ker=CLKernelCreate(prg,"memtests"); if(ker!=INVALID_HANDLE){ int objects=2560; int global_ids();ArrayResize(global_ids,objects,0); ArrayFill(global_ids,0,objects,0); int local_ids();ArrayResize(local_ids,objects,0); ArrayFill(local_ids,0,objects,0); int group_ids();ArrayResize(group_ids,objects,0); ArrayFill(group_ids,0,objects,0); int global_id_handle=CLBufferCreate(ctx,objects*4,CL_MEM_WRITE_ONLY); int local_id_handle=CLBufferCreate(ctx,objects*4,CL_MEM_WRITE_ONLY); int group_id_handle=CLBufferCreate(ctx,objects*4,CL_MEM_WRITE_ONLY); CLSetKernelArgMem(ker,0,global_id_handle); CLSetKernelArgMem(ker,1,local_id_handle); CLSetKernelArgMem(ker,2,group_id_handle); uint offsets()={0}; uint works()={objects}; CLExecute(ker,1,offsets,works); whereas(CLExecutionStatus(ker)!=CL_COMPLETE){Sleep(10);} Print("Kernel completed"); CLBufferRead(global_id_handle,global_ids,0,0,objects); CLBufferRead(local_id_handle,local_ids,0,0,objects); CLBufferRead(group_id_handle,group_ids,0,0,objects); int f=FileOpen("OCLlog.txt",FILE_WRITE|FILE_TXT); for(int i=0;i<objects;i++){ FileWriteString(f,"GLOBAL.ID("+IntegerToString(i)+")="+IntegerToString(global_ids(i))+" : LOCAL.ID("+IntegerToString(i)+")="+IntegerToString(local_ids(i))+" : GROUP.ID("+IntegerToString(i)+")="+IntegerToString(group_ids(i))+"n"); } FileClose(f); int kernel_local_mem_size=CLGetInfoInteger(ker,CL_KERNEL_LOCAL_MEM_SIZE); int kernel_private_mem_size=CLGetInfoInteger(ker,CL_KERNEL_PRIVATE_MEM_SIZE); int kernel_work_group_size=CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE); Print("Kernel native mem ("+kernel_local_mem_size+")"); Print("Kernel personal mem ("+kernel_private_mem_size+")"); Print("Kernel work group dimension ("+kernel_work_group_size+")"); CLKernelFree(ker); CLBufferFree(global_id_handle); CLBufferFree(local_id_handle); CLBufferFree(group_id_handle); }else{Print("Can't create kernel");} CLProgramFree(prg); }else{Alert(errors);} CLContextFree(ctx); } else{ Print("Can't create ctx"); } } void OnTick() { }
Hervorgehoben sind der Artikelparameter und der Export der erhaltenen Indizes in eine Datei.
Auch dieser Code druckt 3 Werte:
- Lokale Speichergröße des Kernels (ich bekomme 1)
- Größe des privaten Kernelspeichers (ich bekomme 0)
- Größe der Kernel-Arbeitsgruppe (ich bekomme 256)
Beginnen wir additionally damit, ihm 100 Gegenstände zuzuwerfen.
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 .. GLOBAL.ID(99)=99 : LOCAL.ID(99)=99 : GROUP.ID(99)=0
So sieht die Dateiausgabe aus (oben) und Sie können sehen, dass die Arbeitslast in keiner Weise aufgeteilt wurde.
Das erste, was mir in den Sinn kommt, ist: « Hat es einen internen Entscheidungsmechanismus? » sowohl für den Wrapper als auch für die nativen mql5-Befehle, wenn sie OpenCL-Befehle aufrufen können. Wie werden wir es wissen? Lassen Sie uns den Kernel etwas « schwerer » machen und auch einen Float-Wert extrahieren.
Lassen Sie uns einen Wertebereich von -2,6 bis 2,6 erstellen und den Tanh dieser x-mal berechnen und auch das Array an Ort und Stelle ändern.
Additionally fügen wir den Kernel-Argumenten einen globalen Float-Zeiger hinzu
__global float* _tangent,
Dann erstellen wir in unserem Programm ein doppeltes Array und füllen es mit zufälligen Werten im Bereich von -2,6 bis 2,6
float tangents();ArrayResize(tangents,objects,0); float vary=5.2; for(int i=0;i<ArraySize(tangents);i++){ float r=(((float)MathRand())/((float)32767.0)*vary)-2.6; tangents(i)=r; }
Erstellen Sie dann ein Puffer-Deal with für openCL, in diesem Fall wird der Speicher gelesen und nicht nur geschrieben. (Lesen und Schreiben für das Gerät, nicht für den Host (uns))
int tangents_handle=CLBufferCreate(ctx,objects*4,CL_MEM_READ_WRITE);
Und verknüpfen Sie dann den Puffer mit den Kernel-Argumenten
CLSetKernelArgMem(ker,3,tangents_handle);
uuund (ich habe hier 2 Stunden verschwendet, weil ich vergessen habe, den Puffer zurückzulesen ????) Vergessen Sie additionally nicht, Buffer Learn aufzurufen, wenn Sie die Daten möchten.
CLBufferRead(tangents_handle,tangents,0,0,objects);
uuund fügen Sie am Ende einen freien Puffer hinzu, da es sonst so aussieht, als würden die Werte (clever) zwischengespeichert
CLBufferFree(tangents_handle);
Cool , jetzt müssen wir es etwas schwieriger machen, indem wir eine Iterationszahl als Argument angeben.
Dadurch wird die tanh-Berechnung wiederholt und für jedes Ergebnis werden wir den tanh der Tangentenfloat / Iterationen summieren.
Sauber.
In diesem Fall müssen wir ein Argument im Kernel hinzufügen, das Argument mit dem Kernel verknüpfen und den Wert festlegen:
"int iterations){" int iterations=100; CLSetKernelArg(ker,4,iterations);
Und die Kernel-Berechnungen, lassen Sie uns den Index von OpenCL nachschlagen
"float sum=(float)0.0;" "float of=(float)_tangent(get_global_id(0));" "for(int i=0;i<iterations;i++){" "sum+=((float)tanh(of-sum))/((float)iterations);" "}" "_tangent(get_global_id(0))=sum;"
- Summe auf 0,0 setzen;
- Legen Sie den Anfangswert in der privaten Speichervariable « of » fest
- Schleife zu Iterationen
- Summieren Sie den Tanh des « von » minus der Summe, die bisher durch die Iterationen geteilt wurde.
- Füllen Sie das Tangentenfeld auf
Einfach und lassen Sie uns die anfänglichen Iterationen auf 100 setzen und den Code erneut ausführen, um zu sehen, ob er immer noch 1 Gruppe erzeugt. (und auch das Ergebnis der Summe zum Debuggen drucken)
Dies ist die exportierte Datei:
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=0.7702 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=0.0282 GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : T(2)=0.9934 GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : T(3)=2.2652 GLOBAL.ID(4)=4 : LOCAL.ID(4)=4 : GROUP.ID(4)=0 : T(4)=-2.2026 ... GLOBAL.ID(96)=96 : LOCAL.ID(96)=96 : GROUP.ID(96)=0 : T(96)=-1.7437 GLOBAL.ID(97)=97 : LOCAL.ID(97)=97 : GROUP.ID(97)=0 : T(97)=-1.1011 GLOBAL.ID(98)=98 : LOCAL.ID(98)=98 : GROUP.ID(98)=0 : T(98)=0.4125 GLOBAL.ID(99)=99 : LOCAL.ID(99)=99 : GROUP.ID(99)=0 : T(99)=1.8560
Okay, es funktioniert, denke ich, und wir bekommen immer noch keine Trennung. Lassen Sie uns die Iterationen auf 10000 erhöhen, aber die Elemente bleiben 100.
Immer noch eine Gruppe
GLOBAL.ID(99)=99 : LOCAL.ID(99)=99 : GROUP.ID(99)=0 : T(99)=0.0905
Lass uns verrückt werden, 10 Millionen Iterationen 100 Elemente, die für die Float-Präzision problematisch sein können? mal sehen
Nein, immer noch eine Gruppe, obwohl die GPU etwas verzögert hat.
GLOBAL.ID(99)=99 : LOCAL.ID(99)=99 : GROUP.ID(99)=0 : T(99)=-2.4797
(Es ist sinnvoll, nicht weiter in Gruppen aufzuteilen, da die Iterationsschleife ein großer Haufen von Berechnungen ist, die keinen Speicher hin und her benötigen, sodass dies optimum für die Ausführung in einem Verarbeitungselement ist, aber ich glaube auch nicht, dass es sich darin aufteilen kann Das Arbeitselement und der Kernel sind eine Arbeitselementinstanz, wenn ich mich nicht irre.
Cool , jetzt lass uns die Iterationen auf 1000 zurücksetzen und mit dem Testen mit unterschiedlichen Artikelmengen beginnen.
Lassen Sie uns 200 Gegenstände darauf werfen
Immer noch eine Gruppe
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=1.5756 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=-1.1957 GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : T(2)=0.6411 ... GLOBAL.ID(198)=198 : LOCAL.ID(198)=198 : GROUP.ID(198)=0 : T(198)=0.5839 GLOBAL.ID(199)=199 : LOCAL.ID(199)=199 : GROUP.ID(199)=0 : T(199)=-1.5742
Jetzt werfen wir 257 Gegenstände!
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=2.0035 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=-0.0069 GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : T(2)=-0.8145 GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : T(3)=1.7880 ... GLOBAL.ID(255)=255 : LOCAL.ID(255)=255 : GROUP.ID(255)=0 : T(255)=0.2042 GLOBAL.ID(256)=256 : LOCAL.ID(256)=256 : GROUP.ID(256)=0 : T(256)=1.7910
Immer noch keine Trennung.
Wie wäre es mit 258? Wir bekommen unseren ersten Cut up!
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=-1.2919 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=-1.2212 ... GLOBAL.ID(127)=127 : LOCAL.ID(127)=127 : GROUP.ID(127)=0 : T(127)=-1.2222 GLOBAL.ID(128)=128 : LOCAL.ID(128)=128 : GROUP.ID(128)=0 : T(128)=1.9752 GLOBAL.ID(129)=129 : LOCAL.ID(129)=0 : GROUP.ID(129)=1 : T(129)=1.0197 ... GLOBAL.ID(255)=255 : LOCAL.ID(255)=126 : GROUP.ID(255)=1 : T(255)=1.9462 GLOBAL.ID(256)=256 : LOCAL.ID(256)=127 : GROUP.ID(256)=1 : T(256)=-1.9560 GLOBAL.ID(257)=257 : LOCAL.ID(257)=128 : GROUP.ID(257)=1 : T(257)=-0.9829
Es teilt die Arbeitsbelastung in zwei Gruppen auf!
Okay, additionally müssen wir das Kriterium für die Aufteilung finden, werfen wir einige Zahlen wie 512 , 1024 , 2048 usw. und protokollieren die Ergebnisse.
512 : Teilt sich wieder in 2 Gruppen auf :
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=-0.3564 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=2.3337 ... GLOBAL.ID(255)=255 : LOCAL.ID(255)=255 : GROUP.ID(255)=0 : T(255)=-2.4480 GLOBAL.ID(256)=256 : LOCAL.ID(256)=0 : GROUP.ID(256)=1 : T(256)=2.3620 ... GLOBAL.ID(510)=510 : LOCAL.ID(510)=254 : GROUP.ID(510)=1 : T(510)=-2.2709 GLOBAL.ID(511)=511 : LOCAL.ID(511)=255 : GROUP.ID(511)=1 : T(511)=-0.3056
Beachten Sie auch, wie praktisch die lokale ID ist, da sie uns den Index in jeder Gruppe mitteilt!
1024: aha, es hat sich jetzt in 4 Gruppen aufgeteilt! Beträgt die maximale Gruppengröße für dieses Gerät additionally 256 Elemente?
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=-0.7910 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : T(1)=-0.7287 ... GLOBAL.ID(255)=255 : LOCAL.ID(255)=255 : GROUP.ID(255)=0 : T(255)=0.2203 GLOBAL.ID(256)=256 : LOCAL.ID(256)=0 : GROUP.ID(256)=1 : T(256)=1.4999 .. GLOBAL.ID(511)=511 : LOCAL.ID(511)=255 : GROUP.ID(511)=1 : T(511)=0.1762 GLOBAL.ID(512)=512 : LOCAL.ID(512)=0 : GROUP.ID(512)=2 : T(512)=-0.0072 ... GLOBAL.ID(767)=767 : LOCAL.ID(767)=255 : GROUP.ID(767)=2 : T(767)=-2.0688 GLOBAL.ID(768)=768 : LOCAL.ID(768)=0 : GROUP.ID(768)=3 : T(768)=-2.0622 ... GLOBAL.ID(1022)=1022 : LOCAL.ID(1022)=254 : GROUP.ID(1022)=3 : T(1022)=2.2044 GLOBAL.ID(1023)=1023 : LOCAL.ID(1023)=255 : GROUP.ID(1023)=3 : T(1023)=-0.6644
Geben wir ihm einen Wert von 768, magazine es gleichmäßige Gruppenbeträge? (wegen seiner Kernzahl?)
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : T(0)=1.8908 ... GLOBAL.ID(255)=255 : LOCAL.ID(255)=255 : GROUP.ID(255)=0 : T(255)=0.0147 GLOBAL.ID(256)=256 : LOCAL.ID(256)=0 : GROUP.ID(256)=1 : T(256)=-1.5271 ... GLOBAL.ID(511)=511 : LOCAL.ID(511)=255 : GROUP.ID(511)=1 : T(511)=2.3339 GLOBAL.ID(512)=512 : LOCAL.ID(512)=0 : GROUP.ID(512)=2 : T(512)=-0.8512 ... GLOBAL.ID(767)=767 : LOCAL.ID(767)=255 : GROUP.ID(767)=2 : T(767)=-0.1783
Nein, keine Probleme, es gibt 3 Gruppen mit jeweils 256 Artikeln. Okay
Können wir hier zu irgendwelchen Schlussfolgerungen kommen? ist der Wert:
int kernel_work_group_size=CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE);
Uns über die maximalen Arbeitsaufgaben zu informieren, die eine Arbeitsgruppe haben kann, denn das macht die GPU selbst, wenn sie keine Anweisungen hat?
Und wenn das stimmt, was ist mit diesem Wert hier?
int device_max_work_group_size=CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE);
Lassen Sie mich es tatsächlich in den Code einfügen und sehen, was es zurückgibt:
1024 steht da. 1024 Arbeitsgruppen oder 1024 Arbeitseinheiten in Arbeitsgruppen insgesamt ?
Nun, finden wir es heraus.
Wenn wir 1025 Arbeitsgruppen (für dieses Gerät) erhalten würden, bräuchten wir 1025*256 Elemente, das sind 262400 Elemente.
Mal sehen . Das ist eine große A**-Datei … aber zum Glück brauchen wir nur die letzte Zeile.
Nun, wir bekommen 1025 Arbeitsgruppen … okay
GLOBAL.ID(262399)=262399 : LOCAL.ID(262399)=255 : GROUP.ID(262399)=1024 : T(262399)=-0.1899
Gehen wir das falsch an? es müssen nicht 1025 parallel ausgeführte Arbeitsgruppen bedeuten, oder?
Was ist, wenn dies anzeigt, dass 1024 Arbeitsgruppen gleichzeitig arbeiten können, weil dieses Gerät 1024 Recheneinheiten hat?
Obwohl es 1 Compute Unit meldet (ich werde diese Data-Befehle auch hinzufügen, damit Sie vergleichen können) und obwohl das cuda-Toolkit 192 cuda-Kerne und 32 Warps meldet.
Das fühlt sich an wie:
- Nvidia hat es halbwegs geschafft
- OpenCL hat es halbwegs geschafft
- Mql5 hat es halbwegs geschafft
und wir halten die Granate und versuchen herauszufinden, ob es eine Weintraube ist oder nicht.
Aber wir mögen keine halben Dinge, wir gehen voll oder gar nicht!
Additionally, wie können wir sagen, was zum Teufel los ist?
Wir müssen die Zeit für die Ausführung messen, wovon aber?
Wir müssen die erhebliche Unterbrechung der Ausführungszeit finden, die darauf hinweist, dass die Kerne Gruppen austauschen, und um die Dinge noch komplizierter zu machen, muss dies vielseitig sein Aktivität im Gerät.
Nicht mein Gerät oder Ihr Gerät, sondern allgemein. (wenn möglich)
Beachten wir auch, dass es kein anderes Ergebnis gibt, als 1/2 * a ** -ing, wenn ein Hersteller nicht vollständig kooperiert. Additionally, mq und khronos sind hier nicht schuld, um truthful zu sein. In diesem Sinne muss ich mich auch auf 0,5 * a ** vorbereiten. ????
Okay, wir müssen jetzt leider eine riesige Klammer öffnen und etwas anderes tun.
Der aktuelle Benchmark hat insofern ein Drawback, als er viel Speicher verbraucht.
Wenn wir wollen, dass der « Cutoff » der Kerne angezeigt wird, müssen wir mehr « Berechnungen » als « Abrufen » verwenden.
Dieser Take a look at wird auch ausgeführt, wenn wir möchten, dass er ständig ausgeführt wird. Es gibt unser zweites Drawback, wenn der Cutoff ist
In der Nähe unserer Schleife, die neu startet, werden wir es nicht bemerken!
Jetzt denken Sie, was ich denke: « Warum lösen Sie das nicht auf den unteren Ebenen, NVIDIA oder AMD »? Und die Antwort ist wahrscheinlich « wir haben nicht all dieses F+E-Geld für Khronos ausgegeben, um herauszukommen und davon zu profitieren » , oder , um es zu vereinfachen « benutze Cuda oder , benutze Hpi » wahrscheinlich.
Weniger schimpfen mehr codieren , oder weniger tippen mehr tippen , außerdem macht meckern nur reich , wenn man anwalt oder politiker ist … ????
Lass es uns angehen, auch wenn es scheitert
Benchmark 2: Ausführungszeitbegrenzung mit Gruppengröße
Stellen wir uns vor, wie können wir die benötigte Zeit messen?
MQ hat eine Möglichkeit bereitgestellt, einen Kernel auszuführen, sobald ich denke . Lassen Sie uns ganz schnell bestätigen.
Ja, wir können auch Größe 0 senden und mit dem Offset spielen, um den GPU-Cache zu umgehen.
Cool additionally, der Benchmark wird:
- Laden Sie einen « schweren » Calcs-Kernel
- Erstellen Sie einen großen Take a look at
- Senden Sie die Elemente einzeln asynchron ???….
hmmmmm, da ist noch ein Drawback. Wir wollen den « Engpass » der GPU (oder des Geräts) finden, aber OpenCL lässt uns das nicht zu, da es die Final selbst handhabt und wir nichts sehen, additionally wie viele Kernel können wir erstellen?
Kann der Benchmark 1000 Kernel sein, deren Ausführung gleichzeitig aufgerufen wird (auch in einer Schleife), und dann beginnen wir mit dem Empfang
Zeit, die ein Kernel benötigt, um fertig zu werden? klingt vernünftig . Um additionally open cl zu testen, bauen wir ein open cl . Eine winzige offene cl ????
Was brauchen wir additionally für den Benchmark?
- Kernels-Array
- Kernel-Startzeit
- Endzeit des Kernels
Dann geben wir die Zeiten aus und entscheiden, wie es weitergeht
Okay, schreiben wir es und machen auch einen ersten Take a look at mit 5 Kerneln gleichzeitig und mit unterschiedlichen Daten!
Hier ist das erste, was wir codieren. Können wir mehrere Instanzen desselben Kernels ausführen? …
class kernel_info{ public: int offset; int deal with; ulong start_microSeconds; ulong end_microSeconds; kernel_info(void){reset();} ~kernel_info(void){reset();} void reset(){ offset=-1; deal with=INVALID_HANDLE; start_microSeconds=0; end_microSeconds=0; } void setup(int _hndl,ulong _start,int _offset){ deal with=_hndl; start_microSeconds=_start; offset=_offset; } void cease(ulong _end){ end_microSeconds=_end; } }; kernel_info KERNELS(); int OnInit() { EventSetMillisecondTimer(33); return(INIT_SUCCEEDED); } void OnDeinit(const int cause) { } void OnTimer(){ EventKillTimer(); int ctx=CLContextCreate(CL_USE_GPU_DOUBLE_ONLY); if(ctx!=INVALID_HANDLE){ string kernel="__kernel void bench(__global int* global_id," "__global int* local_id," "__global int* group_id," "__global float* _tangent," "int iterations){" "float sum=(float)0.0;" "float of=(float)_tangent(get_global_id(0));" "for(int i=0;i<iterations;i++){" "sum+=((float)tanh(of-sum))/((float)iterations);" "}" "sum=(float)0.12345;" "_tangent(get_global_id(0))=0.123;" "global_id(get_global_id(0))=get_global_id(0);" "local_id(get_global_id(0))=get_local_id(0);" "group_id(get_global_id(0))=get_group_id(0);}"; string errors=""; int prg=CLProgramCreate(ctx,kernel,errors); if(prg!=INVALID_HANDLE){ ResetLastError(); int kernels_to_deploy=5; int iterations=1000; ArrayResize(KERNELS,kernels_to_deploy,0); bool deployed=true; for(int i=0;i<kernels_to_deploy;i++){ KERNELS(i).deal with=CLKernelCreate(prg,"bench"); if(KERNELS(i).deal with==INVALID_HANDLE){deployed=false;} } if(deployed){ Print("Deployed all kernels!"); for(int i=0;i<kernels_to_deploy;i++){ if(KERNELS(i).deal with!=INVALID_HANDLE){Print("Kernel ("+i+") Legitimate");} else{Print("Kernel ("+i+") InValid");} } }else{ Print("Can't deploy all kernels!"); for(int i=0;i<kernels_to_deploy;i++){ if(KERNELS(i).deal with!=INVALID_HANDLE){Print("Kernel ("+i+") Legitimate");} else{Print("Kernel ("+i+") InValid");} } } for(int i=0;i<kernels_to_deploy;i++){ if(KERNELS(i).deal with!=INVALID_HANDLE){ CLKernelFree(KERNELS(i).deal with); } } CLProgramFree(prg); }else{Alert(errors);} CLContextFree(ctx); } else{ Print("Can't create ctx"); } }
Nun, natürlich können wir das, dafür ist das hier ????
Obwohl ich eine kleine Verzögerung feststelle.
Lassen Sie uns dies auf bis zu 50 Kernel aufstocken und die Zeit zwischen dem Begin und dem Ende des Timers messen. Ohne etwas anderes zu tun, mounten Sie einfach 50 Kernel auf OpenCL.
Wir verschieben die Kernelmenge außerhalb der if-Blöcke, setzen 2 Timer auf Begin und Ende und geben die Differenz aus:
lengthy timer_ended=GetTickCount(); lengthy diff=timer_ended-timer_started; if(timer_ended<timer_started){diff=UINT_MAX-timer_started+timer_ended;} Print("Time to load and unload "+IntegerToString(kernels_to_deploy)+" kernels = "+IntegerToString(diff)+"ms"); ExpertRemove();
und wir führen es erneut für 5 Kernel aus:
2023.05.02 20:11:41.352 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5 kernels = 94ms
kay jetzt tun 50 Kernel
2023.05.02 20:12:15.704 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 50 kernels = 78ms
genial, mache jetzt 50 Kernel mit einer toten Final (kein Cache)
2023.05.02 20:13:16.359 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 50 kernels = 94ms
Ausgezeichnet, keine Verzögerungen. Lassen Sie uns Massenoperationen mit 5000 Kerneln durchführen. bei 78ms bekommen wir grünes Licht für den Benchmark.
(Es kann jedoch zu einer Druckverzögerung kommen 😀 Drucken Sie additionally nur das aus, was nicht bereitgestellt wird! und den Standing)
2023.05.02 20:15:35.724 blog_kernel_times_benchmark (USDJPY,H1) Deployed all kernels! 2023.05.02 20:15:35.746 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5000 kernels = 94ms
Großartig . Na dann los !
Wir müssen jetzt Berechnungen im Kernel umwandeln in: Wert erhalten, wahnsinnig viele Berechnungen durchführen, Wert angeben.
Wir haben eine Bedingung, die für die Iterationen erfüllt sein muss:
Die Millisekunden, die für die Ausführung EINER Arbeitseinheit benötigt werden, müssen um eine bestimmte Größenordnung größer sein als das Timer-Intervall, damit wir es messen können!
Wenn additionally der « Benchmark » die optimalen Iterationen selbst finden würde, würde er in eine Schleife gehen, bis die gesendeten Iterationen « mal » größer als das Intervall sind.
Lassen Sie mich wiederholen, was wir hier tun: Wir wollen den „Shift“-Wechsel der Verarbeitungselemente „erfassen“, d Elemente (ich vermute, es wird gleich sein
CL_DEVICE_MAX_WORK_GROUP_SIZE
Was ?
Lassen Sie mich erklären, was meiner Meinung nach hier passieren kann:
Wir haben eine Recheneinheit mit 10 Verarbeitungs-Subkernen (die Verarbeitungselemente)
Wenn ich 10 Kernel gleichzeitig zur Ausführung sende, erhalte ich insgesamt eine Ausführungszeit von 150 Millisekunden, was bedeutet, dass meine minimal aufgezeichnete Zeit innerhalb der Kernel-Data-Elemente, die von der maximal aufgezeichneten Zeit abgezogen werden, 150 Millisekunden beträgt.
Wenn ich jetzt 11 Kernel ausführe und ~300 Millisekunden finde, sagt mir das etwas Richtiges, vs. wenn ich 170 Millisekunden gefunden habe.
Rechts ? Ich könnte falsch liegen .
Mal sehen ! Das macht Spaß ! ????