OpenCL: sfide reali - pagina 8

[Roffild]

Mathemat:

Guardate il vostro codice: E poi, nell'ultima riga, voi stessi dividete 240 per 18 (che sono unità per la vostra carta).

È evidente che sei confuso su qualcosa. Ecco il pezzo controverso:

   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]);

Conclusione: globale=30 locale=1

E 240 byte è esattamente quando si crea il buffer.

[Sceptic Philozoff]

Roffild:

È evidente che sei confuso su qualcosa. Ecco un pezzo controverso:

Uscita: globale=30 locale=1

E 240 byte esatti quando si crea il buffer.

globale = 240. Stampa

global_work_size[0]

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

P.S. Sì, hai capito bene. Pardon. Mi sono un po' confuso.

local_work_size[0] = (uint) 30/18 = 1. E io ho lo stesso, dato che unità=28.

[Sceptic Philozoff]

Di nuovo, Roffild:

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

Un compito ben parcellizzato (vedi gli script di MetaDriver dal primo thread ocl) può raggiungere speedup fino a 1000 o più su GPU (rispetto all'esecuzione a 1-thread su CPU su MQL5). Se non riesci a trovarlo - posso mandartelo, puoi testarlo sulla tua carta.

[Roffild]

Avete capito il buffer e la sua velocità?

Faresti meglio a usare AMD CodeXL per capire le UNITA' ecc... - ha dei bei grafici di performance.

AMD CodeXL stesso è glitchato, ma è difficile trarre conclusioni senza di esso.

Non ho intenzione di usare OpenCL finché il tester non mi permette di usare la CPU o finché non ho un compito che dura più a lungo di Number of_buffers * 0.353 msec.

P.S.

Ho finalmente finito di ottimizzare il mio codice e la variante finale supera il test in 33 secondi (320 secondi - prima dell'ottimizzazione, 55 secondi - "OpenCL-style").

[Sceptic Philozoff]

Roffild: Avete capito il buffer e la sua velocità?

Non c'è niente da capire. È chiaro che è un'operazione lenta. Conclusione - aumentate il lavoro all'interno del kernel (ce n'è troppo poco nel vostro codice).

E comprate una scheda video più moderna, sembra che sia diventata migliore con essa.

AMD CodeXL stesso è glitchato, ma è difficile trarre conclusioni senza di esso.

Anche l'utilità di Intel è piuttosto utile, ma per le pietre Intel. Bene, e per catturare gli errori più ovvi nel kernel.

P.S. Ho, dopo tutto, finito di ottimizzare il mio codice e la variante finale ora passa il test in 33 secondi (320 secondi - prima dell'ottimizzazione, 55 secondi - "OpenCL-style").

È già molto meglio.

[Roffild]

Oggi avevo bisogno di generare un array con 1 bit in numeri.

Allo stesso tempo ho fatto pratica con OpenCL.

Sto postando il codice come dimostrazione di un metodo interessante per calcolare global_work_size e local_work_size. L'idea stessa è presa da IntrotoOpenCL.pdf (ne ho una copia), ma l'ho modificata.

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);
}