OpenCL: Echte Herausforderungen - Seite 8

[Roffild]

Mathemat:

Sehen Sie sich Ihren eigenen Code an: Und dann, in der letzten Zeile, dividieren Sie selbst 240 durch 18 (das sind die Einheiten für Ihre Karte).

Sie sind offensichtlich etwas verwirrt. Hier ist der umstrittene Artikel:

   double price[30];

   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); /// <-- здесь же НЕ sizeof(price) вообще-то!
   local_work_size[0] = global_work_size[0] / units;

Print("global=",global_work_size[0]," local=",local_work_size[0]);

Schlussfolgerung: global=30 lokal=1

Und 240 Bytes sind genau dann, wenn Sie den Puffer erstellen.

[Sceptic Philozoff]

Roffild:

Sie sind offensichtlich etwas verwirrt. Hier ist ein kontroverser Artikel:

Ausgabe: global=30 lokal=1

Und genau 240 Bytes bei der Erstellung des Puffers.

global = 240. ausdrucken

global_work_size[0]

Und local_work_size[0] = (uint) 240/18 = 13

P.S. Ja, Sie haben es richtig verstanden. Verzeihung. Ich war ein wenig verwirrt.

local_work_size[0] = (uint) 30/18 = 1. Und ich habe das Gleiche, da Einheiten=28.

[Sceptic Philozoff]

Nochmals, Roffild:

Mathemat: Давай тупо прикинем. 18 задач, выполняемых одновременно на мухах GPU, - это максимум то, что можно сделать на 4-5 нитках CPU. А CPU на x86 эмуляции может организовать гораздо больше ниток. Во всяком случае, если это Intel. Мой бывший Pentium G840 (2 ядра) дал ускорение примерно в 70 раз - на двух unit'ах! Я уже не говорю о том, что вытворяет мой текущий... условно говоря, i7.

Eine gut parallelisierte Aufgabe (siehe die Skripte von MetaDriver aus dem ersten ocl-Thread) kann auf der GPU Geschwindigkeitssteigerungen von bis zu 1000 oder mehr erreichen (im Vergleich zur Ausführung mit einem Thread auf der CPU bei MQL5). Wenn Sie es nicht finden können - ich kann es Ihnen schicken, Sie können es auf Ihrer Karte testen.

[Roffild]

Haben Sie den Puffer und seine Geschwindigkeit herausgefunden?

Verwenden Sie besser AMD CodeXL, um UNITS usw. herauszufinden - es hat schöne Leistungsdiagramme.

AMD CodeXL selbst ist fehlerhaft, aber ohne ihn ist es schwierig, irgendwelche Schlussfolgerungen zu ziehen.

Ich werde OpenCL erst verwenden, wenn der Tester mir erlaubt, die CPU zu verwenden, oder wenn ich eine Aufgabe habe, die länger als Number of_buffers * 0,353 msec dauert.

P.S.

Ich habe die Optimierung meines Codes abgeschlossen und die endgültige Variante besteht den Test in 33 Sekunden (320 Sekunden - vor der Optimierung, 55 Sekunden - "OpenCL-Stil").

[Sceptic Philozoff]

Roffild: Haben Sie den Puffer und seine Geschwindigkeit herausgefunden?

Da gibt es nichts herauszufinden. Es ist klar, dass dies ein langsamer Vorgang ist. Fazit: Mehr Arbeit im Kernel (davon gibt es in Ihrem Code zu wenig).

Und kaufen Sie eine modernere Grafikkarte, es scheint damit besser geworden zu sein.

AMD CodeXL selbst ist fehlerhaft, aber es ist schwer, ohne ihn irgendwelche Schlüsse zu ziehen.

Das Dienstprogramm von Intel ist ebenfalls recht nützlich, allerdings für Intel-Steine. Nun, und um die offensichtlichsten Fehler im Kernel zu erkennen.

P.S. Ich habe schließlich die Optimierung meines Codes abgeschlossen und die endgültige Variante besteht den Test in 33 Sekunden (320 Sekunden - vor der Optimierung, 55 Sekunden - "OpenCL-Stil").

Es ist schon viel besser.

[Roffild]

Heute musste ich ein Array mit 1 Bit in Zahlen erzeugen.

Gleichzeitig habe ich mit OpenCL geübt.

Ich poste den Code als Demonstration einer interessanten Methode zur Berechnung von global_work_size und local_work_size. Die Idee selbst stammt aus IntrotoOpenCL.pdf (ich habe ein Exemplar), aber ich habe sie abgeändert.

void OnStart()
{
   const string source =
      "kernel void differ(const int sizearray, const int bits, global uchar *result)        \r\n"
      "{                                                                                    \r\n"
      "   size_t global_index = get_global_id(0);                                           \r\n"
      "   if (global_index >= sizearray) return; // проверка границ, когда work>arraysize   \r\n"
      "   size_t xor = global_index;                                                        \r\n"
      "   uchar secc = 0;                                                                   \r\n"
      "   for (int bit = bits; bit>-1; bit--)                                               \r\n"
      "     if ((xor & ((size_t)1 << bit)) > 0) secc++;                                     \r\n"
      "   result[global_index] = secc;                                                      \r\n"
      "}                                                                                    \r\n"
   ;
   
   int hContext = CLContextCreate();
   string build_log = "";
   int hProgram = CLProgramCreate(hContext, source, build_log);
   Print("Error = ",build_log);
   int hKernel = CLKernelCreate(hProgram, "differ");
   
   uchar alldiff[1 << 17] = {0};
   CLSetKernelArg(hKernel, 0, ArraySize(alldiff));
   CLSetKernelArg(hKernel, 1, 17 /*bits*/);
   int hBuffer = CLBufferCreate(hContext, sizeof(alldiff), CL_MEM_WRITE_ONLY);
   CLSetKernelArgMem(hKernel, 2, hBuffer);
   CLBufferWrite(hBuffer, alldiff);
   
   /*uchar group_size[1024] = {0};
   uint deviceinfo_size = 0;
   CLGetDeviceInfo(hContext, CL_DEVICE_MAX_WORK_GROUP_SIZE, group_size, deviceinfo_size);
      for (int x = deviceinfo_size; x>=0; x--) Print(group_size[x]);
      Print("ch ",CharArrayToString(group_size));
   */ ///// CLGetDeviceInfo возвращает массив битов (шо за бред?)
   uint group_size = 256;
   
   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(alldiff);
   local_work_size[0] = global_work_size[0] / units;
   if (local_work_size[0] < 1) local_work_size[0] = 1;
   if (local_work_size[0] > group_size) local_work_size[0] = group_size;
   if (global_work_size[0] % local_work_size[0] != 0)
   {
      // увеличиваем global, чтобы global % local == 0
      // в самом kernel проверяется выход за границы
      global_work_size[0] = (int(global_work_size[0] / local_work_size[0]) +1) * local_work_size[0];
      // объяснение в
      // http://wiki.rac.manchester.ac.uk/community/OpenCL?action=AttachFile&amp;do=get&target=IntrotoOpenCL.pdf
   }
      Print("work=", global_work_size[0], " local=", local_work_size[0], " group=", group_size);
   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, alldiff);
   
   int hDump = FileOpen("alldiff.diff", FILE_ANSI|FILE_WRITE);
   for (int x = 0, xcount = ArraySize(alldiff); x < xcount; x++)
      FileWriteString(hDump, (string)alldiff[x]+",");
   FileClose(hDump);
   
   CLBufferFree(hBuffer);
   CLKernelFree(hKernel);
   CLProgramFree(hProgram);
   CLContextFree(hContext);
}