OpenCL: Echte Herausforderungen - Seite 5

[Ilyas]

Erstellen Sie eine Anfrage an den SD und fügen Sie den Code bei (Sie können dies per PM tun), ich werde Ihren Code analysieren.

[Roffild]

An welchem Teil meines Codes sind Sie genau interessiert? Ich habe eine Menge Abhängigkeiten von verschiedenen Dateien.

Das Problem, das ich jetzt habe, ist nur das Schreiben und Lesen des Puffers in 1 Tick des Testers, und für die Überprüfung ist es genug:

#property copyright ""
#property link      ""

int hcontext, hprogram, hkernel, hbuffer[5];

void InitGlobal()
{
   for (int cdev = (int)CLGetInfoInteger(0, CL_DEVICE_COUNT)-1; cdev>-1; cdev--)
   {
      string name;
      CLGetInfoString(cdev, CL_DEVICE_NAME, name);
      Print("Device #",cdev," = ",name);
   }
   
   string source =
"kernel void tester(global double *price, global double *result)                                   \r\n"
"{                                                                                                 \r\n"
"   int global_index = get_global_id(0);                                                           \r\n"
"   result[global_index] = price[global_index] / global_index;                                     \r\n"
"}                                                                                                 \r\n"
;
   
   hcontext = CLContextCreate(CL_USE_GPU_ONLY);
   string build_log;
   hprogram = CLProgramCreate(hcontext, source, build_log);
   Print("build log = ", build_log);
   hkernel = CLKernelCreate(hprogram, "tester");
}

void DeinitGlobal()
{
   CLBufferFree(hbuffer[0]);
   CLBufferFree(hbuffer[1]);
   
   CLKernelFree(hkernel);
   CLProgramFree(hprogram);
   CLContextFree(hcontext);
}

int OnInit()
{
   InitGlobal();
   return(0);
}

void OnDeinit(const int reason)
{
   DeinitGlobal();
}

// Скрипт, в отличии от эксперта, можно дебажить на выходных :)
//void OnStart() {  InitGlobal();
void OnTick() {
   double price[30];
   CopyClose(_Symbol,_Period,0,ArraySize(price),price);
   
   static bool firststart = true;
   if (firststart)
   {
      firststart = false;
      uint bufsize = sizeof(price);
      Print("Размер буфера в байтах =",bufsize);
      hbuffer[0] = CLBufferCreate(hcontext, bufsize, CL_MEM_READ_ONLY);
      hbuffer[1] = CLBufferCreate(hcontext, bufsize, CL_MEM_WRITE_ONLY);
      
      CLSetKernelArgMem(hkernel, 0, hbuffer[0]);
      CLSetKernelArgMem(hkernel, 1, hbuffer[1]);
   }
   
   // А вот здесь не хватает clGetMemObjectInfo(buffer, CL_MEM_SIZE) для проверки размера.
   
   CLBufferWrite(hbuffer[0], price);
   
   uint units = (uint)CLGetInfoInteger(hcontext, CL_DEVICE_MAX_COMPUTE_UNITS);
   uint global_work_offset[] = {0};
   uint global_work_size[1];
   uint local_work_size[1];
   global_work_size[0] = ArraySize(price);
   local_work_size[0] = global_work_size[0] / units;
   bool exec = CLExecute(hkernel, 1, global_work_offset, global_work_size, local_work_size); // async
   if (exec == false) Print("Error in ",__FUNCSIG__," CLExecute: ",GetLastError());
   
   CLBufferRead(hbuffer[1], price);
   
   if (MQL5InfoInteger(MQL5_PROGRAM_TYPE) == PROGRAM_SCRIPT) DeinitGlobal();
}

Läuft nach Skript:

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Gerät #1 = AMD Phenom(tm) II X4 925 Prozessor
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Gerät #0 = Cypress
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) OpenCL: GPU-Gerät 'Cypress' ausgewählt
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Puffergröße in Bytes =240

Experte läuft im Tester vom 2013.01.09 bis 2013.10.10 auf M5 mit "OHLC auf M1":

2013.10.30 19:01:44 Core 1 EURUSD,M5: Test von experts\OpenCL_buffer_test.ex5 von 2013.01.09 00:00 bis 2013.10.10 00:00 gestartet
2013.10.30 19:01:44 Kern 1 2013.01.09 00:00:00 Gerät #0 = Cypress
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 OpenCL: GPU-Gerät 'Cypress' ausgewählt
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 build log =
2013.10.30 19:01:44 Kern 1 2013.01.09 00:00:00 Puffergröße in Bytes =240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 Ticks (55953 Balken) innerhalb von 192443 ms erzeugt (Gesamtbalken in der Historie 131439, Gesamtzeit 192521 ms)
2013.10.30 19:04:55 Kern 1 294 Mb Speicher verwendet

Beachten Sie, dass nur 1 Gerät im Tester vorhanden ist.

Wenn

   //CLBufferRead(hbuffer[1], price);

dann

2013.10.30 19:16:00 Core 1 EURUSD,M5: 1108637 Ticks (55953 Balken) innerhalb von 88218 ms erzeugt (Gesamtbalken in der Historie 131439, Gesamtzeit 88297 ms)

[Ilyas]

Lesen ist erforderlich; tatsächlich ist es die Funktion, die auf die Beendigung des Kernels wartet, da CLExecute den Auftrag nur zur Ausführung in die Warteschlange stellt und dann die Kontrolle an das MQL-Programm zurückgibt, ohne auf dessen Beendigung zu warten

[Ilyas]

In diesem Beispiel wird der Vorteil der Verwendung von OpenCL durch den Overhead beim Kopieren von Puffern wieder aufgezehrt.

Wenn es notwendig ist, Berechnungen mit OHLC-Daten durchzuführen, ist es zwingend erforderlich, sparsam zu schreiben, d. h. einen größeren Puffer im Voraus anzulegen und diesen nur zu überschreiben, wenn neue Daten eintreffen, wobei dem Kernel der neue Beginn und die Größe des Puffers mitgeteilt werden.

[Roffild]

Selbst wenn es uns gelingt, die OHLC-Übertragung zu optimieren (wir verwenden CLSetKernelArg, um den letzten Takt zu übertragen), werden wir beim Lesen des Ergebnispuffers abstürzen:

2013.10.31 19:24:13 Core 1 EURUSD,M5: 1108637 Ticks (55953 Balken) innerhalb von 114489 ms erzeugt (Gesamtbalken in der Historie 131439, Gesamtzeit 114598 ms)

(Verschiebung der Zeile mit CLBufferWrite(hbuffer[0], price); unter IF)

[Sceptic Philozoff]

Roffild: Äh... Artikel über Geschwindigkeitssteigerung mit OpenCL auf der GPU erwiesen sich als Märchen, da sie nicht wirklich auf die anstehenden Aufgaben eingingen :(

Und wer hindert Sie daran, das zu tun? Geh und schreibe etwas Reales, das kein Märchen ist. Versuchen Sie aber, ein Beispiel zu finden, damit eine Beschleunigung stattfindet. Dies ist der schwierigste Teil.

Wenn Sie von meinen Artikeln sprechen... Nun, ich habe eine Fibel geschrieben. Und die Matrixmultiplikation ist eine recht nützliche Operation.

P.S. Übrigens, wenn Ihre CPU von Intel ist, ist die Emulation von x86-Kernen auf ihr viel schneller als auf einer Konkurrenz-CPU. Das heißt, wenn Sie es pro Kern neu berechnen.

HD5850: im Grunde eine ziemlich anständige Karte, aber moderne Karten sind besser - nicht nur aufgrund von mehr Fliegen, sondern auch aufgrund von OpenCL-Optimierungen. So wird beispielsweise die Zugriffszeit auf den globalen Speicher erheblich reduziert.

P.P.S. OpenCL ist kein Allheilmittel; es ist ein brauchbares Werkzeug, das in einigen seltenen Fällen die Geschwindigkeit deutlich erhöhen kann. Und in anderen, nicht so bequemen Fällen ist die Beschleunigung nicht so beeindruckend - wenn es eine gibt.

[Vladimir Gomonov]

Roffild:
Äh... Artikel über Geschwindigkeitssteigerungen mit OpenCL auf GPU haben sich als Märchen herausgestellt, da sie sich nicht wirklich mit realen Aufgaben befassen :(

Nein. Das Märchen ist, dass "jeder Algorithmus in OpenCL beschleunigt werden kann". Nicht jeder Algorithmus.

Der erste Thread zu OpenCL beschreibt sogar recht gut die Kriterien, die ein Algorithmus erfüllen muss, um das Potenzial zur Beschleunigung von OpenCL zu haben.

Viel Glück dabei.

[Roffild]

Die Behauptung bezieht sich nicht auf die Berechnungsgeschwindigkeit - es gibt einen 2-fachen Geschwindigkeitszuwachs (0,02 ms gegenüber 0,05 ms).

Die Behauptung lautet, dass die Artikel keine Informationen enthalten:

• Lese-/Schreiblatenz selbst eines kleinen Puffers = 0,35300 ms - das macht die Umstellung der meisten Algorithmen auf OpenCL zunichte!
• Der Tester wählt keine CPU für OpenCL aus - dies wird nirgends gemeldet!

Ich bin wahrscheinlich der erste, der den Test auf Kosten der GPU beschleunigen wollte, nachdem ich die Versprechen gelesen hatte...

MetaDriver: Der erste Thread zu OpenCL beschreibt sogar recht gut die Kriterien, die ein Algorithmus erfüllen muss, um das Potenzial zur Beschleunigung von OpenCL zu haben.

Lesen Sie meinen Beitrag noch einmal.

Das Hauptkriterium: die Ausführung von MQL-Code im "OpenCL-Stil" sollte die Zeit = Number of_Buffers * 0,35300 ms in 1 Tick überschreiten.

Um die Geschwindigkeit des Algorithmus in MQL mit einer Genauigkeit von Mikrosekunden (1000 Mikrosekunden = 1 Millisekunde) zu ermitteln, müssen Sie ihn mehrmals im Tester und Total_Time / Number_of_Ticks (mein oberster Beitrag) ausführen.

Wäre die Pufferverzögerung nicht, würde mein Code den Test in ~30 Sekunden bestehen - das ist ~2 mal schneller als MQL im "OpenCL-Stil" (55 Sekunden) und ~11 mal schneller als normaler Code (320 Sekunden).

Welche anderen Kriterien gibt es?

[Sceptic Philozoff]

Roffild: Die Behauptung lautet, dass die Artikel keine Informationen enthalten:

• Lese-/Schreiblatenz selbst eines kleinen Puffers = 0,35300 ms - das macht die meisten Algorithmen in OpenCL sinnlos!

Nach Ihrer Erfahrung im Umgang mit OpenCL zu urteilen, haben Sie sicher schon verstanden, dass nicht jeder Algorithmus direkt beschleunigt wird. Eines der Hauptprobleme dabei ist die Minimierung der globalen Speicherzugriffe.

Übrigens muss ich jetzt ein ähnliches Problem mit dem zufälligen Zugriff auf den globalen Gerätespeicher lösen (zu privat dieser zufällige Zugriff, und es ist ein verdammter Overhead). Ich werde das Problem lösen, sobald ich mein Gehirn wieder auf Vordermann gebracht habe.

Der Tester wählt die CPU nicht für OpenCL aus - dies wird nirgends gemeldet!

Schreiben Sie an den Service Desk und begründen Sie den Bedarf an einer solchen Funktion.

Wenn das Prüfgerät nicht verwendet wird, ist es bereits fertig (dies ist meine Anwendung). Und ich habe das Testgerät noch nicht überprüft.

[Roffild]

Mathemat:

Es wurde bereits geschrieben, dass nicht jeder Algorithmus direkt beschleunigt wird. Hier muss man seinen Verstand gebrauchen, und eines der Hauptprobleme besteht darin, die globalen Speicherzugriffe zu minimieren.

Nun, jetzt muss ich ein ähnliches Problem mit dem zufälligen Zugriff auf den globalen Speicher lösen (dieser zufällige Zugriff ist zu häufig). Ich werde das Problem lösen, sobald ich mein Gehirn eingeschaltet habe.

Es ist an der Zeit, Ihr Gehirn zu benutzen, denn 0,35300 ms bezieht sich genau auf clEnqueue[Read/Write]Buffer() und nicht auf globale Speicherzugriffe innerhalb des Kernels.

Das zweite Problem kann durch die Optimierung des Kernels selbst gelöst werden, während das erste eine eiserne Grenze darstellt.