OpenCL: desafíos reales - página 8

[Roffild]

Mathemat:

Mira tu propio código: Y luego, en la última línea, tú mismo divides 240 entre 18 (son unidades para tu tarjeta).

Es evidente que estás confundido en algo. Aquí está la pieza polémica:

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

Conclusión: global=30 local=1

Y 240 bytes es exactamente cuando se crea el buffer.

[Sceptic Philozoff]

Roffild:

Es evidente que estás confundido en algo. Este es un artículo controvertido:

Salida: global=30 local=1

Y 240 bytes exactamente al crear el buffer.

global = 240. Imprimir

global_work_size[0]

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

P.D. Sí, has acertado. Perdón. Me confundí un poco.

local_work_size[0] = (uint) 30/18 = 1. Y yo tengo lo mismo, ya que unidades=28.

[Sceptic Philozoff]

De nuevo, Roffild:

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

Una tarea bien paralelizada (véanse los scripts de MetaDriver del primer hilo ocl) puede lograr aumentos de velocidad de hasta 1000 o más en la GPU (en comparación con la ejecución de un hilo en la CPU en MQL5). Si no lo encuentras, te lo puedo enviar y lo puedes probar en tu tarjeta.

[Roffild]

¿Has averiguado el búfer y su velocidad?

Es mejor que utilices AMD CodeXL para averiguar las UNIDADES, etc. - tiene buenos gráficos de rendimiento.

El propio AMD CodeXL tiene fallos, pero es difícil sacar conclusiones sin él.

No voy a usar OpenCL hasta que el probador me permita usar la CPU o hasta que ejecute una tarea que dure más que Number of_buffers * 0.353 mseg.

P.D.

Terminé de optimizar mi código y la variante final pasa la prueba en 33 segundos (320 segundos - antes de la optimización, 55 segundos - "estilo OpenCL").

[Sceptic Philozoff]

Roffild: ¿Has averiguado el búfer y su velocidad?

No hay nada que averiguar. Está claro que es una operación lenta. Conclusión - aumentar el trabajo dentro del núcleo (hay muy poco en su código).

Y compra una tarjeta de video más moderna, parece que ha mejorado con ella.

El propio AMD CodeXL tiene fallos, pero es difícil sacar conclusiones sin él.

La utilidad de Intel también es bastante útil, pero para las piedras de Intel. Bueno, y para captar los errores más evidentes del núcleo.

P.D. Después de todo, he terminado de optimizar mi código y la variante final pasa la prueba en 33 segundos (320 segundos - antes de la optimización, 55 segundos - "estilo OpenCL").

Ya está mucho mejor.

[Roffild]

Hoy necesitaba generar un array con 1 bit en números.

Al mismo tiempo he practicado con OpenCL.

Pongo el código como demostración de un método interesante para calcular global_work_size y local_work_size. La idea en sí está tomada de IntrotoOpenCL.pdf (tengo una copia), pero la he retocado.

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