OpenCL: desafíos reales - página 5

[Ilyas]

Crea una solicitud a la SD y adjunta el código (puedes hacerlo por MP), analizaré tu código.

[Roffild]

¿Qué parte de mi código le interesa exactamente? Tengo muchas dependencias de diferentes archivos.

El problema que tengo ahora es que sólo escribo y leo el buffer en 1 tick del probador, y para comprobarlo es suficiente:

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

Funcionando según el guión:

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositivo #1 = Procesador AMD Phenom(tm) II X4 925
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositivo #0 = Cypress
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) OpenCL: dispositivo GPU 'Cypress' seleccionado
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Tamaño del buffer en bytes =240

Ejecutando experto en tester desde 2013.01.09 hasta 2013.10.10 en M5 con "OHLC en M1":

2013.10.30 19:01:44 Núcleo 1 EURUSD,M5: se ha iniciado la prueba de experts\OpenCL_buffer_test.ex5 desde 2013.01.09 00:00 hasta 2013.10.10 00:00
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 Dispositivo #0 = Cypress
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 OpenCL: Dispositivo GPU 'Cypress' seleccionado
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 build log =
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 tamaño del buffer en bytes =240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) generados en 192443 ms (total de barras en el historial 131439, tiempo total 192521 ms)
2013.10.30 19:04:55 Núcleo 1 294 Mb de memoria utilizada

Tenga en cuenta que sólo hay un dispositivo en el probador.

Si

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

entonces

2013.10.30 19:16:00 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) generados en 88218 ms (total de barras en el historial 131439, tiempo total 88297 ms)

[Ilyas]

Read es necesaria; de hecho, es la función que espera a que el núcleo termine, ya que CLExecute sólo pone en cola el trabajo para su ejecución y luego devuelve el control al programa MQL sin esperar a que termine

[Ilyas]

En concreto, en este ejemplo, la ventaja de utilizar OpenCL se ve consumida por la sobrecarga de la copia del búfer.

Si es necesario realizar cálculos sobre los datos de OHLC, es imprescindible realizar una escritura ahorradora, es decir, crear un búfer más grande por adelantado y sólo sobrescribir estos nuevos datos cuando lleguen nuevos datos, indicando al núcleo el nuevo comienzo y tamaño del búfer.

[Roffild]

Aunque consigamos optimizar la transferencia de OHLC (utilizaremos CLSetKernelArg para transferir la última barra), seguiremos colapsando al leer el buffer de resultados:

2013.10.31 19:24:13 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) generados en 114489 ms (total de barras en el historial 131439, tiempo total 114598 ms)

(se ha movido la línea con CLBufferWrite(hbuffer[0], price); bajo IF)

[Sceptic Philozoff]

Roffild: Eh... Los artículos sobre el aumento de la velocidad utilizando OpenCL en la GPU resultaron ser un cuento de hadas, ya que no abordaron realmente las tareas en cuestión :(

¿Quién te lo impide? Ve y escribe algo real que no sea un cuento de hadas. Pero trata de encontrar un ejemplo para que la aceleración se produzca. Esta es la parte más difícil.

Si te refieres a mis artículos... Bueno, estaba escribiendo una cartilla. Y la multiplicación de matrices es una operación bastante útil.

P.D. Por cierto, si tu CPU es Intel, la emulación de núcleos x86 en ella es mucho más rápida que en una CPU de la competencia. Eso si se recalcula por núcleo.

HD5850: básicamente es una tarjeta bastante decente, pero las tarjetas modernas son mejores, no sólo porque tienen más moscas, sino también por las optimizaciones de OpenCL. Por ejemplo, el tiempo de acceso a la memoria global se reduce considerablemente.

P.P.S. OpenCL no es una panacea; es una herramienta viable que puede acelerar significativamente en algunos casos raros. Y en otros casos no tan convenientes, la aceleración no es tan impresionante, si es que la hay.

[Vladimir Gomonov]

Roffild:
Eh... Los artículos sobre el aumento de la velocidad utilizando OpenCL en la GPU resultaron ser un cuento de hadas, ya que no se ocupan realmente de las tareas reales :(

No es así. El cuento de hadas es que "cualquier algoritmo puede ser acelerado en OpenCL". No cualquier algoritmo.

El primer hilo sobre OpenCL incluso describe bastante bien los criterios que debe poseer un algoritmo para tener potencial de aceleración ocl.

Buena suerte con eso.

[Roffild]

La afirmación no se refiere a la velocidad de cálculo: hay un aumento de velocidad del doble (0,02 ms frente a 0,05 ms)

La reclamación es que no hay información en los artículos:

• La latencia de lectura/escritura, incluso de un búfer pequeño, es de 0,35300 ms. ¡Esto es lo que invalida la conversión de la mayoría de los algoritmos a OpenCL!
• El probador no selecciona la CPU para OpenCL - ¡esto no se informa en ninguna parte!

Probablemente soy el primero que quería acelerar la prueba a costa de la GPU, después de leer las promesas...

MetaDriver: El primer hilo sobre OpenCL incluso describe bastante bien los criterios que debe poseer un algoritmo para tener potencial de aceleración ocl.

Vuelve a leer mi post.

El criterio principal: la ejecución del código MQL en el "estilo OpenCL" debe superar el tiempo = Número de_Buffers * 0,35300 ms en 1 tick.

Para averiguar la velocidad del algoritmo en MQL con una precisión de microsegundos (1000 microsegundos = 1 milisegundo), tendrá que ejecutar en el probador varias veces y Total_Time / Number_of_Ticks (mi puesto superior).

Si no fuera por el retardo del búfer, mi código pasaría la prueba en ~30 segundos, es decir, ~2 veces más rápido que el MQL "estilo OpenCL" (55 segundos) y ~11 veces más rápido que el código normal (320 segundos).

¿Qué otros criterios hay?

[Sceptic Philozoff]

Roffild: La reclamación es que no hay información en los artículos:

• La latencia de lectura/escritura, incluso de un búfer pequeño, es de 0,35300 ms. ¡Esto es lo que hace que la conversión de la mayoría de los algoritmos a OpenCL no tenga sentido!

A juzgar por su experiencia en el manejo de OpenCL, ya debe haber comprendido que no todos los algoritmos se aceleran directamente. Uno de los principales problemas aquí es minimizar los accesos a la memoria global.

Por cierto, ahora tengo que resolver un problema similar con el acceso aleatorio a la memoria global del dispositivo (demasiado privado este acceso aleatorio, y es una puta sobrecarga). Lo resolveré en cuanto recupere mi cerebro.

El probador no selecciona la CPU para OpenCL - ¡esto no se informa en ninguna parte!

Escriba al Servicio de Atención al Cliente y justifique la necesidad de dicha función.

Si el probador no se utiliza, ya está hecho (esta es mi aplicación). Y todavía no he comprobado el probador.

[Roffild]

Mathemat:

Ya se ha escrito que no todos los algoritmos se aceleran directamente. Aquí hay que usar el cerebro, y uno de los principales problemas es minimizar los accesos a la memoria global.

Pues bien, ahora tengo que resolver un problema similar con el acceso aleatorio a la memoria global (este acceso aleatorio es demasiado frecuente). Lo resolveré en cuanto me funcione el cerebro.

Es hora de usar tu cerebro porque 0,35300 ms se refiere exactamente a clEnqueue[Read/Write]Buffer() y no a los accesos globales a la memoria dentro del kernel.

La segunda puede resolverse optimizando el propio núcleo, mientras que la primera es una limitación de hierro.