OpenCL: desafios reais - página 5

[Ilyas]

Crie um pedido ao SD e anexe o código (pode fazê-lo via PM), eu analisarei o seu código.

[Roffild]

Em que parte do meu código está exactamente interessado? Tenho muitas dependências em ficheiros diferentes.

O problema que tenho agora é apenas escrever e ler o tampão em 1 carrapato do testador, e para o verificar é 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();
}

Executado por guião:

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositivo #1 = Fenómeno AMD(tm) II X4 925 Processador
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositivo #0 = Cipreste
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) Tamanho do buffer em bytes =240

Perito em testes de 2013.01.09 a 2013.10.10 em M5 com "OHLC em M1":

2013.10.30 19:01:44 Core 1 EURUSD,M5: testes de peritos\OpenCL_buffer_test.ex5 de 2013.01.09 00:00 a 2013.10.10 00:00
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 Dispositivo #0 = Cipreste
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 tamanho do buffer em bytes =240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) gerados em 192443 ms (total de barras na história 131439, tempo total 192521 ms)
2013.10.30 19:04:55 Core 1 294 Mb memória utilizada

Note-se que existe apenas 1 dispositivo no aparelho de teste.

Se

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

então

2013.10.30 19:16:00 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) gerados dentro de 88218 ms (total de barras na história 131439, tempo total 88297 ms)

[Ilyas]

É necessário ler; de facto, é a função à espera que o kernel termine, uma vez que a CLExecute apenas enfileira o trabalho para execução e depois devolve o controlo ao programa MQL sem esperar que este termine

[Ilyas]

Especificamente neste exemplo, a vantagem de utilizar o OpenCL é consumida pelas despesas gerais de cópia tampão.

Se for necessário efectuar cálculos sobre dados OHLC, então é imperativo fazer uma escrita parcimoniosa, ou seja, criar um buffer maior com antecedência e só sobrescrever estes novos dados quando novos dados chegarem, dizendo ao kernel o novo começo e tamanho do buffer.

[Roffild]

Mesmo que consigamos optimizar a transferência OHLC (utilizaremos CLSetKernelArg para transferir a última barra), continuaremos a falhar ao ler o buffer de resultados:

2013.10.31 19:24:13 Core 1 EURUSD,M5: 1108637 ticks (55953 barras) gerados dentro de 114489 ms (total de barras na história 131439, tempo total 114598 ms)

(moveu a linha com CLBufferWrite(hbuffer[0], preço); em IF)

[Sceptic Philozoff]

Roffild: Eh.... Artigos sobre o aumento da velocidade usando OpenCL em GPU revelaram-se um conto de fadas, uma vez que não abordavam realmente o objectivo :(

Bem, quem o impede de o fazer? Vai e escreve algo real que não seria um conto de fadas. Mas tente encontrar um exemplo para que a aceleração aconteça. Esta é a parte mais difícil.

Se estás a falar dos meus artigos... ...estava a escrever uma cartilha. E a multiplicação matricial é uma operação bastante útil.

P.S. A propósito, se o seu CPU for Intel, a emulação de núcleos x86 nele é muito mais rápida do que num CPU concorrente. Isto é, se o recalcularmos por núcleo.

HD5850: basicamente um cartão bastante decente, mas os cartões modernos são melhores - não só devido a mais moscas, mas também devido a optimizações OpenCL. Por exemplo, o tempo de acesso à memória global é significativamente reduzido.

P.P.S. OpenCL não é uma panaceia; é uma ferramenta viável que pode acelerar significativamente em alguns casos raros. E em outros casos não tão convenientes, a aceleração não é tão impressionante - se é que existe uma.

[Vladimir Gomonov]

Roffild:
Eh.... Artigos sobre aceleração usando OpenCL em GPU revelaram-se um conto de fadas, uma vez que não lidam realmente com tarefas reais :(

Não é assim. O conto de fadas é que "qualquer algoritmo pode ser acelerado no OpenCL", não qualquer algoritmo.

O primeiro fio do OpenCL descreve até bastante bem os critérios que um algoritmo deve possuir para ter um potencial de aceleração ocl.

Boa sorte com isso.

[Roffild]

A reclamação não se refere à velocidade de cálculo - há uma velocidade 2x (0,02 ms vs 0,05 ms)

A alegação é que não há informação nos artigos:

• Latência de leitura/escrita mesmo de um pequeno tampão = 0,35300 ms - isto é o que invalida a conversão da maioria dos algoritmos para OpenCL!
• O testador não selecciona CPU para OpenCL - isto não é relatado em lado nenhum!

Sou provavelmente o primeiro a querer acelerar o teste à custa da GPU, tendo lido as promessas.

MetaDriver: O primeiro fio do OpenCL descreve até bastante bem os critérios que um algoritmo deve possuir para ter um potencial de aceleração ocl.

Leia novamente o meu post.

O critério principal: a execução do código MQL no "estilo OpenCL" deve exceder o tempo = Número de_Buffers * 0,35300 ms em 1 tick.

Para descobrir a velocidade do algoritmo em MQL com uma precisão de microssegundos (1000 microssegundos = 1 milissegundo), terá de executá-lo várias vezes no testador e no Total_Time / Number_of_Ticks (o meu posto superior).

Se não fosse o atraso do tampão, o meu código passaria o teste em ~30 segundos - isto é ~2 vezes mais rápido que o MQL "estilo OpenCL" (55 segundos) e ~11 vezes mais rápido que o código normal (320 segundos).

Que outros critérios existem?

[Sceptic Philozoff]

Roffild: A alegação é que não há informação nos artigos:

• Latência de leitura/escrita mesmo de um pequeno tampão = 0,35300 ms - isto é o que torna a maioria dos algoritmos de conversão para OpenCL sem sentido!

A julgar pela sua experiência em lidar com o OpenCL, já deve ter compreendido que nem todos os algoritmos são directamente acelerados. Um dos principais problemas aqui é a minimização do acesso à memória global.

A propósito, tenho agora de resolver um problema semelhante com o acesso aleatório à memória global do dispositivo (demasiado privado este acesso aleatório, e é uma porra de uma sobrecarga). Vou resolvê-lo assim que tiver o meu cérebro de volta ao normal.

O testador não selecciona a CPU para OpenCL - isto não é relatado em lado nenhum!

Escrever ao Service Desk e justificar a necessidade de tal recurso.

Se o testador não for utilizado, já está feito (esta é a minha aplicação). E ainda não verifiquei o testador.

[Roffild]

Mathemat:

Já foi escrito que nem todos os algoritmos são directamente acelerados. Aqui tem de usar o seu cérebro, e um dos principais problemas é minimizar o acesso à memória global.

Bem, agora tenho de resolver um problema semelhante com o acesso aleatório à memória global (este acesso aleatório é demasiado frequente). Vou resolvê-lo assim que puser o meu cérebro a trabalhar.

É tempo de usar o cérebro porque 0,35300 ms refere-se exactamente ao clEnqueue[Read/Write]Buffer() e não a acessos globais à memória dentro do núcleo.

A segunda pode ser resolvida através da optimização do próprio núcleo enquanto a primeira é uma limitação de ferro.