OpenCL: desafios reais - página 4

[Vladimir Kazakov]

TheXpert:

O método mais difícil é o FeedPatterns.

Num relance, não há forma de o acelerar: há muito pouco trabalho para os grãos relativamente à quantidade de dados (muitos dados, pouco trabalho), todos os ganhos serão consumidos pela cópia para trás e para a frente.

Pode tentar pedir explicitamente ao compilador que o faça em paralelo em VS ou criar fios paralelos para o processador.

[Vladimir Gomonov]

TheXpert:
Eventualmente, irei transferi-lo para a MQL.

Será estritamente necessário? Há muito tempo que quero investigar o AMP e agora há uma hipótese...

[Mykola Demko]

kazakov.v:

Num relance, não há forma de acelerar: há muito pouco trabalho para os grãos em relação à quantidade de dados (muitos dados, pouco trabalho), todo o ganho será consumido pela cópia para a frente e para trás.

Pode tentar dizer explicitamente ao compilador para o paralelizar em VS ou criar roscas paralelas para o processador.

Porquê, estas tarefas são perfeitamente adequadas ao 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;
  }

E o material verde também o pode fazer.

[Mykola Demko]

Urain:

Porquê, estas tarefas são óptimas no OpenCL

e também se pode fazer o material verde.

Precisamos de fazer um teste de implementação e comparação em OpenCL e C++, e se houver um aumento, então traduzir tudo.

[Roffild]

O CL_DEVICE_PREFERRED_VECTOR_VECTOR_WIDTH_* indica o tamanho máximo do vector ou o tamanho óptimo?

Quando CL_DEVICE_PREFERRED_VECTOR_VECTOR_LARGURA_DOUBLE=2 já será o dobro3 e o dobro4 lento?

[Sceptic Philozoff]

Roffild:

O CL_DEVICE_PREFERRED_VECTOR_VECTOR_WIDTH_* indica o tamanho máximo do vector ou o tamanho óptimo?

Quando CL_DEVICE_PREFERRED_VECTOR_VECTOR_LARGURA_DOUBLE=2 já será o dobro3 e o dobro4 lento?

1. Ao máximo.

2- É pouco provável que a desaceleração seja significativa, mas não irá aumentar a velocidade de execução.

[Roffild]

O MQL OpenCL é apenas um invólucro sobre o API original.
Preciso de saber respostas e esclarecimentos sobre a implementação deste invólucro.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Isto é realmente uma fila para um dispositivo, não um contexto?

Os amortecedores são de leitura/escrita síncronos ou assíncronos?
clEnqueue[Leitura/Escrita]Buffer(consulta, buffer, CL_TRUE) - aqui CL_TRUE ou CL_FALSE ?

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

O CLExecute() devolve imediatamente o controlo? Não o bloqueia durante o tempo da sua execução?
Parece que são precisos 2-40 ms para o colocar na fila.

Agora aqui está a questão principal:
Quando e sob que condições é chamado clFinish()? E por causa da ausência de clFinish(), é difícil fazer fila.

E a ajuda MQL não descreve de todo CL_MEM_*_HOST_PTR, mas eles estão presentes lá.

Finalmente converti completamente o meu indicador para o estilo OpenCL.
Realização do teste de 2013.01.09 a 2013.10.10 em M5 com "OHLC em M1":
320 segundos - antes da tradução
55 segundos - Emulação estilo OpenCL em MQL5:

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

Mas a corrida da GPU foi frustrante para mim :(
Esperava fazer o teste em menos de 30 segundos, mas tive um atraso completo para CLBufferWrite!

Carregamento da placa de vídeo a 32% e aprovação no teste em 1710 segundos sem CL_MEM_*_HOST_PTR
Carregamento da placa de vídeo a 22% e realização de um teste em 740 segundos com CL_MEM_ALLOC_HOST_PTR
CL_MEM_HOST_PTR e CL_MEM_USE_HOST_PTR resultam em CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Então, como proceder ao intercâmbio adequado de dados?

E ainda não é seleccionada nenhuma CPU para cálculos no testador.

Adaptador de vídeo = ATI Radeon HD 5850
Processador = Fenómeno AMD(tm) II X4 925 Processador

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Isto é realmente uma fila para um dispositivo, não um contexto?

Sim, o contexto e a fila são criados por dispositivo (a investigação mostrou que o opencl não funciona correctamente com vários dispositivos diferentes).

Os amortecedores são de leitura/escrita síncronos ou assíncronos?
clEnqueue[Leitura/Escrita]Buffer(pergunta, buffer, CL_TRUE) - aqui CL_TRUE ou CL_FALSE ?

A leitura e a escrita são síncronas.

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - não implementado.
O CLExecute() devolve imediatamente o controlo? Não o bloqueia durante o tempo da sua execução?

Sim

E agora a questão principal:
Quando e sob que condições é clFinish() chamado ? E devido à ausência de clFinish(), é difícil formar uma fila de espera.

Não chamado, lido de memória deve ser utilizado.

E não há qualquer descrição de CL_MEM_*_HOST_PTR na Ajuda MQL.

Finalmente converti completamente o meu indicador para o estilo OpenCL.
Realização do teste de 2013.01.09 a 2013.10.10 em M5 com "OHLC em M1":
320 segundos - antes da tradução
55 segundos - Emulação estilo OpenCL em MQL5:
Mas a corrida da GPU foi frustrante para mim :(
Esperava que o teste funcionasse em menos de 30 ms e tive um atraso total para CLBufferWrite!

Carregamento da placa de vídeo a 32% e aprovação no teste em 1710 segundos sem CL_MEM_*_HOST_PTR
Carregamento da placa de vídeo a 22% e testes em 740 segundos com CL_MEM_ALLOC_HOST_PTR
CL_MEM_HOST_PTR e CL_MEM_USE_HOST_PTR resultam em CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Então, como proceder ao intercâmbio adequado de dados?

As bandeiras CL_MEM_HOST_PTR e CL_MEM_USE_HOST_PTR não são actualmente suportadas pelo terminal (estamos a investigar esta questão).

E ainda não é seleccionada nenhuma CPU para cálculos no testador.

Já tentou especificar explicitamente o dispositivo CPU?

[Roffild]

Que tal tentar dar-nos amortecedores assíncronos e clFinish() ?

Há uma suposição de que é a escrita síncrona que está a abrandar, o que até mesmo o AMD CodeXL indica:

"clEnqueueWriteBuffer: Sincronização desnecessária. Bloqueio da escrita".

E no testador de CPU não é sequer seleccionável por número. Bug #865549.

[Roffild]

Eh.... Artigos sobre aumento de velocidade usando OpenCL em GPU revelaram-se um conto de fadas por não lidarem com tarefas reais :(

Este mês, escrevi milhares de linhas de código para conquistar o OpenCL.

Assim, para depurar o OpenCL, tive de emular funções de MQL para as executar através do AMD CodeXL em C/C+++.

Vou repetir os resultados dos testes de 2013.01.09 a 2013.10.10 em M5 com "OHLC em M1":
320 segundos - antes da tradução
55 segundos - Emulação estilo OpenCL em MQL5

O "estilo OpenCL" é reduzir ao mínimo o número de chamadas CopyHigh/CopyTime/CopyOpen/... e aumentar a quantidade de código para processar arrays depois destas funções serem chamadas.

E estes cálculos são o que faltam nos belos artigos sobre OpenCL:

Resultado do teste sem OpenCL:
Núcleo 1 EURUSD,M5: 1108637 ticks (55953 barras) gerados dentro de 55427 ms (total de barras na história 131439, tempo total 55520 ms)
55427 ms / 1108637 tick = 0,04999 ms/tick - 1 tick por CPU (execução em OpenCL não deve exceder este tempo)

Isto é o que consegui ao executar o meu próprio código em C/C++ e ao executá-lo através do AMD CodeXL:
0,02000 ms - 0,05000 ms - execução do meu kernel na GPU
0,35300 ms - uma chamada para clEnqueueWriteBuffer para 168 bytes a 500KB/s
0,35300 ms - uma chamada clEnqueueWriteBuffer para 3,445 KBytes com 9,500 MBytes/s (o tempo médio de transferência é o mesmo)

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

E recebi 3,445 KByte devido a um erro de cálculo do tamanho do conjunto 21*168, mas mesmo isso não afectou o tempo de transferência.

Conclusão: mesmo se eu conseguir optimizar o meu kernel para 0,02000 ms, o que é de facto ~2 vezes mais rápido que a habitual passagem MQL (0,04999 ms), tudo se resume à velocidade de leitura/escrita da GPU (0,35300 ms - ~7 vezes mais lento que o cálculo MQL!).

O CPU não é seleccionado no meu testador para OpenCL, por isso não posso usar mais 3 núcleos vazios...

P.S.
55 segundos ainda não é o limite de optimização em MQL, é apenas uma emulação OpenCL quando não há suporte :)