OpenCL: Echte Herausforderungen - Seite 4

[Vladimir Kazakov]

TheXpert:

Die schwierigste Methode ist FeedPatterns.

Auf den ersten Blick gibt es keine Möglichkeit, dies zu beschleunigen: Der Arbeitsaufwand für die Kernel ist im Verhältnis zur Datenmenge sehr gering (viele Daten, wenig Arbeit), alle Gewinne werden durch das Hin- und Herkopieren wieder aufgefressen.

Sie können versuchen, den Compiler explizit aufzufordern, das Programm in VS zu parallelisieren oder parallele Threads für den Prozessor zu erstellen.

[Vladimir Gomonov]

TheXpert:
Letztendlich werde ich es auf MQL übertragen.

Ich wollte mich schon lange mit AMP beschäftigen, und jetzt gibt es eine Chance...

[Mykola Demko]

kazakov.v:

Auf den ersten Blick gibt es keine Möglichkeit, dies zu beschleunigen: Der Arbeitsaufwand für den Kernel ist im Verhältnis zur Datenmenge sehr gering (viele Daten, wenig Arbeit), der gesamte Gewinn wird durch das Hin- und Herkopieren wieder aufgezehrt.

Sie können versuchen, dem Compiler explizit mitzuteilen, dass er das Programm in VS parallelisieren soll, oder parallele Threads für den Prozessor erstellen.

Diese Aufgaben sind perfekt für OpenCL geeignet

void Solver::FeedPattern(const Pattern &pattern)
  {
   size_t size=pattern.Size();

   assert(size==m_PatternSize);
   if(size!=m_PatternSize)
     {
      return;
     }

   const std::vector<double>&values=pattern.Values();
   double etalon=pattern.Etalon();

   size_t i;

   for(i=0; i<size;++i)
     {
      for(size_t j=0; j<size;++j)
        {
         m_Matrix[i][j]+=values[i]*values[j];
        }
      m_Matrix[i][size]+=values[i];
      m_Matrix[i][size+1]+=values[i]*etalon;
     }

   for(i=0; i<size;++i)
     {
      m_Matrix[size][i]+=values[i];
     }
   m_Matrix[size][size]+=1;
   m_Matrix[size][size+1]+=etalon;
  }

Und das Grünzeug kann es auch.

[Mykola Demko]

Urain:

Diese Aufgaben sind für OpenCL hervorragend geeignet

und du kannst auch das grüne Zeug machen.

Wir müssen eine Implementierung und einen Vergleichstest in OpenCL und C++ durchführen, und wenn es eine Steigerung gibt, dann alles übersetzen.

[Roffild]

Gibt CL_DEVICE_PREFERRED_VECTOR_WIDTH_* die maximale Vektorgröße oder die optimale Größe an?

Wenn CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 ist, sind dann double3 und double4 bereits langsam?

[Sceptic Philozoff]

Roffild:

Gibt CL_DEVICE_PREFERRED_VECTOR_WIDTH_* die maximale Vektorgröße oder die optimale Größe an?

Wenn CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 ist, sind dann double3 und double4 bereits langsam?

1. Bis zum Maximum.

2- Eine Verlangsamung ist unwahrscheinlich, aber die Ausführungsgeschwindigkeit wird dadurch nicht erhöht.

[Roffild]

MQL OpenCL ist nur ein Wrapper über der ursprünglichen API.
Ich benötige Antworten und Klarstellungen zur Implementierung dieses Wrappers.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Handelt es sich dabei um eine Warteschlange für ein Gerät und nicht um einen Kontext?

Sind die Puffer zum Lesen und Schreiben synchron oder asynchron?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - hier CL_TRUE oder CL_FALSE ?

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - nicht implementiert.

Gibt CLExecute() die Kontrolle sofort zurück? Sie blockiert ihn nicht für die Zeit seiner Ausführung?
Es scheint 2-40 ms zu dauern, bis es in die Warteschlange gestellt wird...

Jetzt kommt die wichtigste Frage:
Wann und unter welchen Bedingungen wird clFinish() aufgerufen? Und weil clFinish() fehlt, ist es schwierig, sich in eine Warteschlange einzureihen.

Und die MQL-Hilfe beschreibt CL_MEM_*_HOST_PTR überhaupt nicht, aber sie sind dort vorhanden.

Ich habe meinen Indikator endlich vollständig auf OpenCL umgestellt.
Durchführung des Tests vom 2013.01.09 bis 2013.10.10 auf M5 mit "OHLC auf M1":
320 Sekunden - vor der Übersetzung
55 Sekunden - Emulation im OpenCL-Stil auf MQL5:

// подготовка данных общая и копия kernel на MQL5, а эмуляция через:
for (int get_global_id = maxcount-1; get_global_id>-1; get_global_id--) NoCL(params,m_result,get_global_id);

Aber der GPU-Lauf war für mich frustrierend :(
Ich hatte gehofft, den Test in weniger als 30 Sekunden durchführen zu können, erhielt aber eine totale Verzögerung für CLBufferWrite!

Laden der Grafikkarte mit 32 % und Bestehen des Tests in 1710 Sekunden ohne CL_MEM_*_HOST_PTR
Laden der Grafikkarte bei 22 % und Durchführung eines Tests in 740 Sekunden mit CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR und CL_MEM_USE_HOST_PTR führen zu CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Wie kann man also Daten richtig austauschen?

Und immer noch ist keine CPU für die Berechnungen im Tester ausgewählt.

Videoadapter = ATI Radeon HD 5850
Prozessor = AMD Phenom(tm) II X4 925-Prozessor

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Handelt es sich dabei um eine Warteschlange für ein Gerät und nicht um einen Kontext?

Ja, Kontext und Warteschlange werden pro Gerät erstellt (Untersuchungen haben gezeigt, dass opencl mit mehreren verschiedenen Geräten nicht korrekt funktioniert).

Sind die Puffer zum Lesen und Schreiben synchron oder asynchron?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - hier CL_TRUE oder CL_FALSE ?

Lesen und Schreiben sind synchron.

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - nicht implementiert.
Gibt CLExecute() die Kontrolle sofort zurück? Sie blockiert ihn nicht für die Zeit seiner Ausführung?

Ja

Und nun die wichtigste Frage:
Wann und unter welchen Bedingungen wird clFinish() aufgerufen? Und wegen des Fehlens von clFinish() ist es schwierig, eine Warteschlange zu bilden.

Nicht aufgerufen, aus dem Speicher lesen muss verwendet werden.

Und in der MQL-Hilfe gibt es überhaupt keine Beschreibung von CL_MEM_*_HOST_PTR.

Ich habe meinen Indikator endlich vollständig auf OpenCL umgestellt.
Durchführung des Tests vom 2013.01.09 bis 2013.10.10 auf M5 mit "OHLC auf M1":
320 Sekunden - vor der Übersetzung
55 Sekunden - Emulation im OpenCL-Stil auf MQL5:
Aber der GPU-Lauf war für mich frustrierend :(
Ich hatte gehofft, dass der Test in weniger als 30 ms abläuft, aber ich bekam einen totalen Lag für CLBufferWrite!

Laden der Grafikkarte mit 32 % und Bestehen des Tests in 1710 Sekunden ohne CL_MEM_*_HOST_PTR
Laden der Grafikkarte bei 22 % und Testen in 740 Sekunden mit CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR und CL_MEM_USE_HOST_PTR führen zu CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Wie kann man also Daten richtig austauschen?

Die Flags CL_MEM_COPY_HOST_PTR und CL_MEM_USE_HOST_PTR werden derzeit vom Terminal nicht unterstützt (wir untersuchen dieses Problem).

Und immer noch ist keine CPU für die Berechnungen im Tester ausgewählt.

Haben Sie versucht, das CPU-Gerät explizit anzugeben?

[Roffild]

Wie wäre es, wenn Sie uns asynchrone Puffer und clFinish() zur Verfügung stellen würden ?

Es besteht die Vermutung, dass das synchrone Schreiben verlangsamt wird, worauf sogar AMD CodeXL hinweist:

"clEnqueueWriteBuffer: Unnötige Synchronisierung. Schreibblockade"

Und im CPU-Tester ist sie nicht einmal nach Nummern auswählbar. Fehler #865549.

[Roffild]

Äh... Artikel über Geschwindigkeitssteigerungen mit OpenCL auf GPUs entpuppten sich als Märchen, da sie sich nicht mit realen Aufgaben befassten :(

Diesen Monat habe ich Tausende von Codezeilen geschrieben, um OpenCL zu erobern.

Um OpenCL zu debuggen, musste ich also Funktionen von MQL emulieren, um sie durch AMD CodeXL in C/C++ laufen zu lassen.

Ich werde die Testergebnisse vom 2013.01.09 bis 2013.10.10 auf M5 mit "OHLC auf M1" wiederholen:
320 Sekunden - vor der Übersetzung
55 Sekunden - Emulation im OpenCL-Stil auf MQL5

Der "OpenCL-Stil" besteht darin, die Anzahl der Aufrufe von CopyHigh/CopyTime/CopyOpen/... auf ein Minimum zu reduzieren. und den Umfang des Codes zur Verarbeitung von Arrays nach dem Aufruf dieser Funktionen zu erhöhen.

Und diese Berechnungen sind es, die in den schönen Artikeln über OpenCL fehlen:

Testergebnis ohne OpenCL:
Core 1 EURUSD,M5: 1108637 Ticks (55953 Balken) generiert innerhalb von 55427 ms (Gesamtbalken in der Historie 131439, Gesamtzeit 55520 ms)
55427 ms / 1108637 Ticks = 0,04999 ms/Tick - 1 Tick pro CPU (die Ausführung unter OpenCL sollte diese Zeit nicht überschreiten)

Das habe ich herausgefunden, als ich meinen eigenen Code in C/C++ durch AMD CodeXL laufen ließ:
0,02000 ms - 0,05000 ms - Ausführung meines Kernels auf der GPU
0,35300 ms - ein Aufruf von clEnqueueWriteBuffer für 168 Byte bei 500KB/s
0,35300 ms - ein clEnqueueWriteBuffer-Aufruf für 3,445 KBytes mit 9,500 MBytes/s (die durchschnittliche Übertragungszeit ist die gleiche)

168 Bytes ist:
double open[21]={1.3668,1.3661,1.36628,1.3664,1.36638,1.36629,1.3664,1.36881,1.36814,1.3692,1.36918,1.36976,1.36816,1.36776,1.36779,1.3695,1.36927,1.36915,1.3679,1.36786,1.36838};

Und wegen eines Berechnungsfehlers bei der Array-Größe von 21*168 erhielt ich 3.445 KByte, aber selbst das hatte keine Auswirkungen auf die Übertragungszeit.

Fazit: Selbst wenn es mir gelingt, meinen Kernel auf 0,02000 ms zu optimieren, was in der Tat ~2 mal schneller ist als der übliche MQL-Durchlauf (0,04999 ms), läuft alles auf die GPU-Lese-/Schreibgeschwindigkeit hinaus (0,35300 ms - ~7 mal langsamer als die MQL-Berechnung!).

CPU ist in meinem Tester nicht für OpenCL ausgewählt, also kann ich keine weiteren 3 leeren Kerne verwenden...

P.S.
55 Sekunden ist noch nicht die Grenze der Optimierung in MQL, es ist nur eine OpenCL-Emulation, wenn es keine Unterstützung gibt :)