In diesem Weblog untersuche ich, wie der lokale Speicher in Bezug auf eine Arbeitsgruppe (von Arbeitselementen) funktioniert.
Wir erstellen einen einfachen Kernel, der IDs, globale ID, lokale ID, Gruppen-ID eines Arbeitselements exportiert.
Zusätzlich werden wir eine lokale Ganzzahl mit dem Präfix __local innerhalb der Kernelfunktion instanziieren
und wir werden es erhöhen (++) .
Der lokale Speicher wird innerhalb einer Arbeitsgruppe geteilt, additionally werden wir sehen, was mit der Ausführung passiert, wenn wir den Wert davon erfassen
-Variable und übergeben Sie sie an den globalen Slot des Arbeitselements, das ausgeführt wird.
Au ja, wir werden auch ein Ausgabearray zusammen mit den IDs haben und den Wert, den es gesehen hat, dort hineinwerfen.
Während wir additionally eine lokale Variable haben, erhalten wir ihren Zustand über mehrere Arbeitselemente hinweg.
das ist der Code:
#property model "1.00" int OnInit() { EventSetMillisecondTimer(33); return(INIT_SUCCEEDED); } void OnDeinit(const int purpose) { } 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 int* output){rn" "//initialized as soon as in native reminiscence for every compute unitrn" "__local int f;" "f++;" "output(get_global_id(0))=f;" "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=8; 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); int output();ArrayResize(output,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); int output_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); CLSetKernelArgMem(ker,3,output_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); CLBufferRead(output_handle,output,0,0,objects); int f=FileOpen("OCLlocalmemtestlog.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))+" : OUTPUT("+IntegerToString(i)+")="+IntegerToString(output(i))+"n"); } FileClose(f); int groups_created=group_ids(0); for(int i=0;i<ArraySize(group_ids);i++){ if(group_ids(i)>groups_created){groups_created=group_ids(i);} } int compute_units=CLGetInfoInteger(ker,CL_DEVICE_MAX_COMPUTE_UNITS); 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); int device_max_work_group_size=CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE); Print("Kernel native mem ("+kernel_local_mem_size+")"); Print("Kernel non-public mem ("+kernel_private_mem_size+")"); Print("Kernel work group measurement ("+kernel_work_group_size+")"); Print("Gadget max work group measurement("+device_max_work_group_size+")"); Print("Gadget max compute items("+compute_units+")"); Print("Gadget Native Mem Dimension ("+CLGetInfoInteger(ctx,CL_DEVICE_LOCAL_MEM_SIZE)+")"); Print("------------------"); Print("Teams created : "+IntegerToString(groups_created+1)); CLKernelFree(ker); CLBufferFree(global_id_handle); CLBufferFree(local_id_handle); CLBufferFree(group_id_handle); CLBufferFree(output_handle); }else{Print("Can not create kernel");} CLProgramFree(prg); }else{Alert(errors);} CLContextFree(ctx); } else{ Print("Can not create ctx"); } }
Wir geben den Inhalt der id – Arrays und des Output – Arrays in eine Datei aus und inspizieren sie . Lassen Sie es uns für 8 Elemente ausführen, wir müssen 256 Elemente in diesem Gerät überschreiten, damit es automatisch in Gruppen aufgeteilt wird, wie von gemeldet Größe der Kernel-Arbeitsgruppe. ,Wir haben 8 Artikel und daher wird es nur 1 Gruppe geben.
Hier ist die Dateiausgabe:
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : OUTPUT(0)=1 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : OUTPUT(1)=1 GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : OUTPUT(2)=1 GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : OUTPUT(3)=1 GLOBAL.ID(4)=4 : LOCAL.ID(4)=4 : GROUP.ID(4)=0 : OUTPUT(4)=1 GLOBAL.ID(5)=5 : LOCAL.ID(5)=5 : GROUP.ID(5)=0 : OUTPUT(5)=1 GLOBAL.ID(6)=6 : LOCAL.ID(6)=6 : GROUP.ID(6)=0 : OUTPUT(6)=1 GLOBAL.ID(7)=7 : LOCAL.ID(7)=7 : GROUP.ID(7)=0 : OUTPUT(7)=1
Wir können alle gedruckten IDs und die Ausgabe sehen, alle Werte sind 1.
Sie haben wahrscheinlich erwartet, dass, da alle diese Elemente gleichzeitig ausgeführt werden, der Anfangswert, den sie vor dem ++ sahen, 0 battle.
Die Spezifikationen besagen, dass wir die lokale Ganzzahl f für die Recheneinheit instanziieren, aber eine Arbeitsgruppe in einer Recheneinheit ausgeführt wird. Die nächste Frage ist additionally, ob wir sie auch professional Arbeitsgruppe instanziieren?
Finden wir es heraus, fügen wir ein native() uint hinzu, um die Ausführungsfunktion zu senden, wie in den vorherigen Blogs zu sehen, um die Arbeit in 2 Arbeitsgruppen mit jeweils 4 Arbeitselementen aufzuteilen.
Wir werden wahrscheinlich die gleiche Ausgabe sehen und die einzige Änderung wird in den lokalen IDs und Gruppen-IDs sein
Dies ist die Zeile, die wir über der Ausführungsfunktion hinzufügen, und um sie zu verwenden, fügen wir sie einfach als letztes Argument in der Ausführungsfunktion hinzu.
(in dieser Dimension geben wir 4 Artikel professional Gruppe an)
uint native()={4}; CLExecute(ker,1,offsets,works,native);
Wie erwartet erstellt unser Code 2 Gruppen:
2023.05.04 00:59:05.922 blog_simple_local_mem_operation (USDJPY,H1) Teams created : 2
Und das ist die Ausgabedatei:
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : OUTPUT(0)=1 GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : OUTPUT(1)=1 GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : OUTPUT(2)=1 GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : OUTPUT(3)=1 GLOBAL.ID(4)=4 : LOCAL.ID(4)=0 : GROUP.ID(4)=1 : OUTPUT(4)=1 GLOBAL.ID(5)=5 : LOCAL.ID(5)=1 : GROUP.ID(5)=1 : OUTPUT(5)=1 GLOBAL.ID(6)=6 : LOCAL.ID(6)=2 : GROUP.ID(6)=1 : OUTPUT(6)=1 GLOBAL.ID(7)=7 : LOCAL.ID(7)=3 : GROUP.ID(7)=1 : OUTPUT(7)=1
Wie erwartet wird die lokale Speicher-Ganzzahl f professional Arbeitsgruppe instanziiert (oder zugewiesen? wie Sie es nennen).
Eindrucksvoll.
Aber was ist, wenn Sie möchten, dass der Wert (von f) innerhalb des Arbeitselements erhöht wird, um es zu verwenden?
Es gibt Befehle, um das mit dem zu tun atomar_ Präfix, in diesem Fall interessieren wir uns für die atomic_inc .
Was sie tun, ist im Wesentlichen, den Bereich um die Variable f zu « bewachen », bis die Arbeitsaufgabe sie ändert, additionally denke ich, dass sie einen kleinen Einbruch in die Geschwindigkeit hat.
(Ich hoffe, ich schlachte hier nicht die Erklärung ab)
Lassen Sie uns additionally eine Model des Obigen schreiben, die sowohl den atomaren als auch den nicht atomaren Wert exportiert, wir werden diese entsprechend benennen
der code sieht jetzt so aus:
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 int* atomic_output," "__global int* non_atomic_output){rn" "//initialized as soon as in native reminiscence for every compute unitrn" "__local int with_atomic,without_atomic;" "with_atomic=0;" "without_atomic=0;" "atomic_inc(&with_atomic);" "without_atomic++;" "atomic_output(get_global_id(0))=with_atomic;" "non_atomic_output(get_global_id(0))=without_atomic;" "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=8; 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); int atomic_output();ArrayResize(atomic_output,objects,0); int non_atomic_output();ArrayResize(non_atomic_output,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); int atomic_output_handle=CLBufferCreate(ctx,objects*4,CL_MEM_WRITE_ONLY); int non_atomic_output_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); CLSetKernelArgMem(ker,3,atomic_output_handle); CLSetKernelArgMem(ker,4,non_atomic_output_handle); uint offsets()={0}; uint works()={objects}; uint native()={4}; CLExecute(ker,1,offsets,works,native); 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); CLBufferRead(atomic_output_handle,atomic_output,0,0,objects); CLBufferRead(non_atomic_output_handle,non_atomic_output,0,0,objects); int f=FileOpen("OCLlocalmemtestlog.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))+" : ATOMIC.OUTPUT("+IntegerToString(i)+")="+IntegerToString(atomic_output(i))+" : NON-ATOMIC.OUTPUT("+IntegerToString(i)+")="+IntegerToString(non_atomic_output(i))+"n"); } FileClose(f); int groups_created=group_ids(0); for(int i=0;i<ArraySize(group_ids);i++){ if(group_ids(i)>groups_created){groups_created=group_ids(i);} } int compute_units=CLGetInfoInteger(ker,CL_DEVICE_MAX_COMPUTE_UNITS); 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); int device_max_work_group_size=CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE); Print("Kernel native mem ("+kernel_local_mem_size+")"); Print("Kernel non-public mem ("+kernel_private_mem_size+")"); Print("Kernel work group measurement ("+kernel_work_group_size+")"); Print("Gadget max work group measurement("+device_max_work_group_size+")"); Print("Gadget max compute items("+compute_units+")"); Print("Gadget Native Mem Dimension ("+CLGetInfoInteger(ctx,CL_DEVICE_LOCAL_MEM_SIZE)+")"); Print("------------------"); Print("Teams created : "+IntegerToString(groups_created+1)); CLKernelFree(ker); CLBufferFree(global_id_handle); CLBufferFree(local_id_handle); CLBufferFree(group_id_handle); CLBufferFree(atomic_output_handle); CLBufferFree(non_atomic_output_handle); }else{Print("Can not create kernel");} CLProgramFree(prg); }else{Alert(errors);} CLContextFree(ctx); } else{ Print("Can not create ctx"); } }
Additionally initialisieren wir professional Gruppe die Variablen with_atomic und without_atomic
und wir werden auch ihre Werte exportieren. Lassen Sie es uns mit denselben Elementen und lokalen Elementen ausführen
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : ATOMIC.OUTPUT(0)=4 : NON-ATOMIC.OUTPUT(0)=1
GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : ATOMIC.OUTPUT(1)=4 : NON-ATOMIC.OUTPUT(1)=1
GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : ATOMIC.OUTPUT(2)=4 : NON-ATOMIC.OUTPUT(2)=1
GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : ATOMIC.OUTPUT(3)=4 : NON-ATOMIC.OUTPUT(3)=1
GLOBAL.ID(4)=4 : LOCAL.ID(4)=0 : GROUP.ID(4)=1 : ATOMIC.OUTPUT(4)=4 : NON-ATOMIC.OUTPUT(4)=1
GLOBAL.ID(5)=5 : LOCAL.ID(5)=1 : GROUP.ID(5)=1 : ATOMIC.OUTPUT(5)=4 : NON-ATOMIC.OUTPUT(5)=1
GLOBAL.ID(6)=6 : LOCAL.ID(6)=2 : GROUP.ID(6)=1 : ATOMIC.OUTPUT(6)=4 : NON-ATOMIC.OUTPUT(6)=1
GLOBAL.ID(7)=7 : LOCAL.ID(7)=3 : GROUP.ID(7)=1 : ATOMIC.OUTPUT(7)=4 : NON-ATOMIC.OUTPUT(7)=1
Uuuund wir bekommen das….hmmm
Das Atom gibt uns den letzten Wert, den es hatte, warum?
Nun, wenn wir uns den Code ansehen, übergeben wir den Wert der lokalen Variablen quick am Ende der Ausführung der Arbeitsgruppe an das globale with_atomic_output-Array.
Additionally stell dir das vor:
- 4 Arbeitselemente (der ersten Gruppe) treten zur Ausführung in die Recheneinheit ein
- Jeder wird einem Verarbeitungselement zugewiesen
- CU initialisiert die 2 Ganzzahlen with_atomic und without_atomic
- Jedes Arbeitselement wird parallel ausgeführt
- Im Allgemeinen ist eine Berechnung viel schneller als eine Übertragung in den globalen Speicher
- Und wir können das anders als atomic_inc(); Funktion nichts anderes hält zurück
die Arbeitselemente, bis jedes den Punkt erreicht, an dem es Daten an das with_atomic-Array zurücksenden soll. - Wenn additionally jedes Ingredient die Ausgangsstufe erreicht, ist der Wert von with_atomic bereits 4.
Wir können es vielleicht sehen, wenn wir ein privates int in den atomaren Aufruf einfügen und 1 hinzufügen.
Die Khronos-Dokumentation gibt dies für die an atomic_inc()
Lesen Sie den 32-Bit-Wert (als alt ) an der Stelle gespeichert, auf die verwiesen wird P . Berechnen ( alt + 1) und speichern Sie das Ergebnis an der Stelle, auf die gezeigt wird P . Die Funktion kehrt zurück alt .
Es sagt uns additionally, dass wir den alten Wert der lokalen Variablen erhalten, wenn wir eine Ganzzahl hyperlinks vom atomaren Aufruf platzieren.
Das bedeutet, dass die lokale Variable gesperrt wird und dann die non-public Ganzzahl den Wert an diesem Punkt erhält, dann wird die Operation (++ weil inc()) für die lokale Variable ausgeführt und dann wird sie entsperrt. Wir erhalten additionally den Wert dieses Arbeitselements « benutzt ».
Dann wird die Variable, die wir haben, privat sein, wir schlagen +1 darauf und erhalten den Wert in dieser Instanz!
Lassen Sie uns das tatsächlich auch exportieren.
Wir fügen einen weiteren Puffer hinzu, ein weiteres Argument, wir verknüpfen den Puffer mit dem Kernel, wir rufen den Wert am Ende ab, wir drucken ihn in die Datei und wir vergessen nicht, den Puffer freizugeben.
Jetzt sieht der Kernel für den Teil, der das « alte » aus dem atomaren Aufruf erhält, so aus:
int this_item_only; this_item_only=atomic_inc(&with_atomic)+1;
Erstellt ein privates int für jedes Ingredient mit dem Namen this_item_only , ruft dann den alten Wert der lokalen Variablen with_atomic ab und fügt ihm eins hinzu
– Entschuldigen Sie meine Variablennamen, es ist für den Take a look at und von mehreren Assessments, die ich ausgeführt habe –
Hier ist die Dateiausgabe:
GLOBAL.ID(0)=0 : LOCAL.ID(0)=0 : GROUP.ID(0)=0 : ATOMIC.OUTPUT(0)=4 : NON-ATOMIC.OUTPUT(0)=1 : INSTANCE.OUTPUT(0)=1
GLOBAL.ID(1)=1 : LOCAL.ID(1)=1 : GROUP.ID(1)=0 : ATOMIC.OUTPUT(1)=4 : NON-ATOMIC.OUTPUT(1)=1 : INSTANCE.OUTPUT(1)=2
GLOBAL.ID(2)=2 : LOCAL.ID(2)=2 : GROUP.ID(2)=0 : ATOMIC.OUTPUT(2)=4 : NON-ATOMIC.OUTPUT(2)=1 : INSTANCE.OUTPUT(2)=3
GLOBAL.ID(3)=3 : LOCAL.ID(3)=3 : GROUP.ID(3)=0 : ATOMIC.OUTPUT(3)=4 : NON-ATOMIC.OUTPUT(3)=1 : INSTANCE.OUTPUT(3)=4
GLOBAL.ID(4)=4 : LOCAL.ID(4)=0 : GROUP.ID(4)=1 : ATOMIC.OUTPUT(4)=4 : NON-ATOMIC.OUTPUT(4)=1 : INSTANCE.OUTPUT(4)=1
GLOBAL.ID(5)=5 : LOCAL.ID(5)=1 : GROUP.ID(5)=1 : ATOMIC.OUTPUT(5)=4 : NON-ATOMIC.OUTPUT(5)=1 : INSTANCE.OUTPUT(5)=2
GLOBAL.ID(6)=6 : LOCAL.ID(6)=2 : GROUP.ID(6)=1 : ATOMIC.OUTPUT(6)=4 : NON-ATOMIC.OUTPUT(6)=1 : INSTANCE.OUTPUT(6)=3
GLOBAL.ID(7)=7 : LOCAL.ID(7)=3 : GROUP.ID(7)=1 : ATOMIC.OUTPUT(7)=4 : NON-ATOMIC.OUTPUT(7)=1 : INSTANCE.OUTPUT(7)=4
Hölle ja.
Als ich dies zum ersten Mal tat, verwendete ich 512 Elemente (statt 8), wodurch ich ein weiteres mögliches Downside finden konnte:
Ich erhielt « ATOMIC.OUTPUT »-Werte von 224 in der zweiten Gruppe statt 256 (Gruppengröße battle 256 professional Gruppe, 2 Gruppen)
Dies wurde durch den Versatz in der Ausführung für einige der Arbeitselemente erreicht, 32 Elemente begannen später, was dazu führte, dass sie das nicht erreichten
atomic_inc-Teil noch und zur gleichen Zeit hatten die anderen 224 Arbeitselemente ihren Export in globale Speicherstufen erreicht, wodurch 224 als atomare Ausgabe gemeldet wurden.
Hier ist ein Schema:
Die 224 Elemente erreichen den Teil des globalen Speichers, wenn die zurückgesetzten 32 Elemente den Teil atomic_inc noch nicht erreicht haben
Sie haben die Lösung für die Artikel gesehen, und zwar die einer Barriere.
Was die Barriere tun wird, ist, alle GRUPPENARTIKEL an der Linie zu stoppen, an der Sie sie platzieren, bis ALLE ANDEREN GRUPPENARTIKEL diese Linie ebenfalls erreichen.
Das hat das Downside gelöst, das ist die Codezeile:
barrier(CLK_GLOBAL_MEM_FENCE);
Wenn wir CLK_LOCAL_MEM_FENCE angeben, dann könnten die Elemente IN DER GRUPPE nichts mit dem lokalen Speicher tun, bis ALLE ELEMENTE dieser Gruppe diese Zeile erreicht haben.
Wo glaubst du, ist diese Linie hingegangen?
Richtig, vor dem Export in den globalen Speicher.
Ich hänge den endgültigen Quellcode an.
Beifall
