OpenCL: desafíos reales - página 4

[Vladimir Kazakov]

TheXpert:

El método más difícil es FeedPatterns.

A simple vista, no hay forma de acelerarlo: hay muy poco trabajo para los núcleos en relación con la cantidad de datos (muchos datos, poco trabajo), todo lo que se gane se lo comerá la copia de ida y vuelta.

Puedes intentar pedir explícitamente al compilador que lo paralelice en VS o crear hilos paralelos para el procesador.

[Vladimir Gomonov]

TheXpert:
Con el tiempo, voy a transferirlo a MQL.

¿Es estrictamente necesario? Llevo mucho tiempo queriendo investigar sobre AMP, y ahora hay una posibilidad...

[Mykola Demko]

kazakov.v:

A simple vista, no hay forma de acelerarlo: hay muy poco trabajo para los núcleos en relación con la cantidad de datos (muchos datos, poco trabajo), toda la ganancia se la comerá la copia de ida y vuelta.

Puedes intentar decirle explícitamente al compilador que lo paralelice en VS o crear hilos paralelos para el procesador.

Por eso, estas tareas se adaptan perfectamente a OpenCL

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

Y lo verde también puede hacerlo.

[Mykola Demko]

Urain:

Por qué, estas tareas son geniales en OpenCL

y también puedes hacer las cosas verdes.

Tenemos que hacer una prueba de implementación y comparación en OpenCL y C++, y si hay un aumento, entonces traducir todo.

[Roffild]

¿Indica CL_DEVICE_PREFERRED_VECTOR_WIDTH_* el tamaño máximo del vector o el tamaño óptimo?

Cuando CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, ¿el double3 y el double4 serán ya lentos?

[Sceptic Philozoff]

Roffild:

¿Indica CL_DEVICE_PREFERRED_VECTOR_WIDTH_* el tamaño máximo del vector o el tamaño óptimo?

Cuando CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, ¿el double3 y el double4 serán ya lentos?

1. Al máximo.

2- La ralentización es poco probable que sea significativa, pero no aumentará la velocidad de ejecución.

[Roffild]

MQL OpenCL es sólo una envoltura sobre la API original.
Necesito saber respuestas y aclaraciones sobre la implementación de este wrapper.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
¿Se trata en realidad de una cola para un dispositivo, no de un contexto?

¿Los búferes de lectura/escritura son síncronos o asíncronos?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - aquí CL_TRUE o CL_FALSE ?

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

¿Devuelve CLExecute() el control inmediatamente? ¿No lo bloquea durante el tiempo de su ejecución?
Parece que tarda entre 2 y 40 ms en ponerse en la cola...

Esta es la cuestión principal:
¿Cuándo y bajo qué condiciones se llama a clFinish()? Y debido a la ausencia de clFinish(), es difícil hacer cola.

Y la ayuda de MQL no describe CL_MEM_*_HOST_PTR en absoluto, pero están presentes allí.

Finalmente he convertido completamente mi indicador al estilo OpenCL.
Realizando la prueba desde 2013.01.09 hasta 2013.10.10 en M5 con "OHLC en M1":
320 segundos - antes de la traducción
55 segundos - Emulación estilo OpenCL en MQL5:

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

Pero la carrera de la GPU fue frustrante para mí :(
¡Esperaba realizar la prueba en menos de 30 segundos pero recibí un retraso total para CLBufferWrite!

Carga de la tarjeta de vídeo al 32% y paso de la prueba en 1710 segundos sin CL_MEM_*_HOST_PTR
Cargando la tarjeta de vídeo al 22% y haciendo una prueba en 740 segundos con CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR y CL_MEM_USE_HOST_PTR resultan en CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Entonces, ¿cómo intercambiar correctamente los datos?

Y todavía no se selecciona ninguna CPU para los cálculos en el probador.

Adaptador de vídeo = ATI Radeon HD 5850
Procesador = Procesador AMD Phenom(tm) II X4 925

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
¿Se trata en realidad de una cola para un dispositivo, no de un contexto?

Sí, el contexto y la cola se crean por dispositivo (la investigación ha demostrado que opencl no funciona correctamente con varios dispositivos diferentes).

¿Los búferes de lectura/escritura son síncronos o asíncronos?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - aquí CL_TRUE o CL_FALSE ?

La lectura y la escritura son sincrónicas.

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - no implementado.
¿Devuelve CLExecute() el control inmediatamente? ¿No lo bloquea durante el tiempo de su ejecución?

Sí

Y ahora la pregunta principal:
¿Cuándo y bajo qué condiciones se llama a clFinish()? Y debido a la ausencia de clFinish(), es difícil formar una cola.

No se llama, se debe utilizar la lectura de la memoria.

Y no hay ninguna descripción de CL_MEM_*_HOST_PTR en la ayuda de MQL.

Finalmente he convertido completamente mi indicador al estilo OpenCL.
Realizando la prueba desde 2013.01.09 hasta 2013.10.10 en M5 con "OHLC en M1":
320 segundos - antes de la traducción
55 segundos - Emulación estilo OpenCL en MQL5:
Pero la carrera de la GPU fue frustrante para mí :(
¡Esperaba que la prueba se ejecutara en menos de 30 ms y obtuve un retraso total para CLBufferWrite!

Carga de la tarjeta de vídeo al 32% y paso de la prueba en 1710 segundos sin CL_MEM_*_HOST_PTR
Cargando la tarjeta de vídeo al 22% y probando en 740 segundos con CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR y CL_MEM_USE_HOST_PTR resultan en CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Entonces, ¿cómo intercambiar correctamente los datos?

Las banderas CL_MEM_COPY_HOST_PTR y CL_MEM_USE_HOST_PTR no son soportadas actualmente por el terminal (estamos investigando este problema).

Y todavía no se selecciona ninguna CPU para los cálculos en el probador.

¿Has probado a especificar explícitamente el dispositivo de la CPU?

[Roffild]

¿Qué tal si intentamos darnos búferes asíncronos y clFinish() ?

Se supone que lo que se ralentiza es la escritura sincrónica, algo que incluso AMD CodeXL insinúa:

"clEnqueueWriteBuffer: Sincronización innecesaria. Bloqueo de la escritura"

Y en el probador de CPU ni siquiera se puede seleccionar por número. Bug #865549.

[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 trataban de tareas reales :(

Este mes he escrito miles de líneas de código para conquistar OpenCL.

Así, para depurar OpenCL, tuve que emular funciones de MQL para ejecutarlas a través de AMD CodeXL en C/C++.

Repetiré los resultados de las pruebas desde el 2013.01.09 hasta el 2013.10.10 en M5 con "OHLC en M1":
320 segundos - antes de la traducción
55 segundos - Emulación estilo OpenCL en MQL5

El "estilo OpenCL" consiste en reducir el número de llamadas CopyHigh/CopyTime/CopyOpen/... al mínimo. y aumentando la cantidad de código para procesar las matrices después de llamar a estas funciones.

Y estos cálculos son los que faltan en los bonitos artículos sobre OpenCL:

Resultado de la prueba sin OpenCL:
Core 1 EURUSD,M5: 1108637 ticks (55953 barras) generados en 55427 ms (total de barras en el historial 131439, tiempo total 55520 ms)
55427 ms / 1108637 tick = 0,04999 ms/tick - 1 tick por CPU (la ejecución en OpenCL no debería superar este tiempo)

Esto es lo que he obtenido ejecutando mi propio código en C/C++ y pasándolo por AMD CodeXL:
0,02000 ms - 0,05000 ms - ejecución de mi kernel en la GPU
0,35300 ms - una llamada a clEnqueueWriteBuffer para 168 bytes a 500KB/s
0,35300 ms - una llamada a clEnqueueWriteBuffer para 3,445 KBytes con 9,500 MBytes/s (el tiempo medio de transferencia es el mismo)

168 Bytes es:
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};

Y obtuve 3.445 KByte debido a un error de cálculo del tamaño de la matriz 21*168, pero incluso eso no afectó al tiempo de transferencia.

Conclusión: aunque consiga optimizar mi kernel a 0,02000 ms, lo que es efectivamente ~2 veces más rápido que el paso habitual de MQL (0,04999 ms), todo se reduce a la velocidad de lectura/escritura de la GPU (0,35300 ms - ¡7 veces más lento que el cálculo de MQL!).

La CPU no está seleccionada en mi probador para OpenCL, así que no puedo usar otros 3 núcleos vacíos...

P.D.
55 segundos no es todavía el límite de optimización en MQL, es sólo una emulación de OpenCL cuando no hay soporte :)