OpenCL: sfide reali - pagina 5

[Ilyas]

Crea una richiesta alla SD e allega il codice (puoi farlo via PM), analizzerò il tuo codice.

[Roffild]

Quale parte del mio codice ti interessa esattamente? Ho un sacco di dipendenze da diversi file.

Il problema che ho ora è solo scrivere e leggere il buffer in 1 tick del tester, e per il controllo è sufficiente:

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

Esecuzione di uno script:

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositivo #1 = Processore 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' selezionato
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dimensione del buffer in byte =240

Esecuzione di expert in tester dal 2013.01.09 al 2013.10.10 su M5 con "OHLC su M1":

2013.10.30 19:01:44 Core 1 EURUSD,M5: test di esperti\OpenCL_buffer_test.ex5 dal 2013.01.09 00:00 al 2013.10.10 00:00 iniziato
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' selezionato
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 Dimensione del buffer in byte =240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 ticks (55953 barre) generati entro 192443 ms (totale barre nella storia 131439, tempo totale 192521 ms)
2013.10.30 19:04:55 Core 1 294 Mb di memoria utilizzata

Si noti che c'è solo 1 dispositivo nel tester.

Se

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

poi

2013.10.30 19:16:00 Core 1 EURUSD,M5: 1108637 ticks (55953 barre) generati entro 88218 ms (totale barre nella storia 131439, tempo totale 88297 ms)

[Ilyas]

Read è richiesto; infatti, è la funzione che aspetta che il kernel termini, poiché CLExecute accoda solo il lavoro per l'esecuzione e poi restituisce il controllo al programma MQL senza aspettare che termini

[Ilyas]

In particolare in questo esempio, il vantaggio di usare OpenCL è consumato dall'overhead della copia del buffer.

Se è necessario eseguire calcoli sui dati OHLC, allora è imperativo fare una scrittura parsimoniosa, cioè creare un buffer più grande in anticipo e sovrascrivere questi nuovi dati solo quando arrivano nuovi dati, dicendo al kernel il nuovo inizio e la dimensione del buffer.

[Roffild]

Anche se riusciamo a ottimizzare il trasferimento OHLC (useremo CLSetKernelArg per trasferire l'ultima barra), andremo comunque in crash durante la lettura del buffer dei risultati:

2013.10.31 19:24:13 Core 1 EURUSD,M5: 1108637 ticks (55953 barre) generati entro 114489 ms (totale barre nella storia 131439, tempo totale 114598 ms)

(spostato la linea con CLBufferWrite(hbuffer[0], prezzo); sotto IF)

[Sceptic Philozoff]

Roffild: Eh... Gli articoli sull'aumento della velocità usando OpenCL su GPU si sono rivelati una favola in quanto non hanno realmente affrontato i compiti da svolgere :(

Ebbene, chi vi impedisce di farlo? Vai a scrivere qualcosa di reale che non sia una favola. Ma cercate di trovare un esempio in modo che l'accelerazione avvenga. Questa è la parte più difficile.

Se stai parlando dei miei articoli... Beh, stavo scrivendo una guida. E la moltiplicazione delle matrici è un'operazione abbastanza utile.

P.S. A proposito, se la vostra CPU è Intel, l'emulazione dei core x86 su di essa è molto più veloce che su una CPU concorrente. Questo se si ricalcola per core.

HD5850: fondamentalmente una scheda abbastanza decente, ma le schede moderne sono migliori - non solo a causa di più mosche, ma anche per le ottimizzazioni OpenCL. Per esempio, il tempo di accesso alla memoria globale è significativamente ridotto.

P.P.S. OpenCL non è una panacea; è uno strumento valido che può accelerare significativamente in alcuni rari casi. E in altri casi non così convenienti, l'accelerazione non è così impressionante - se c'è.

[Vladimir Gomonov]

Roffild:
Eh... Gli articoli sulla velocizzazione usando OpenCL su GPU si sono rivelati una favola, dato che non si occupano realmente di compiti reali :(

Non è così. La favola è che "qualsiasi algoritmo può essere accelerato in OpenCL". Non qualsiasi algoritmo.

Il primo thread su OpenCL descrive anche abbastanza bene i criteri che un algoritmo deve possedere per avere un potenziale di accelerazione ocl.

Buona fortuna.

[Roffild]

L'affermazione non riguarda la velocità di calcolo - c'è un aumento di velocità di 2 volte (0,02 ms contro 0,05 ms)

L'affermazione è che non ci sono informazioni negli articoli:

• Latenza di lettura/scrittura anche di un piccolo buffer = 0,35300 ms - questo è ciò che invalida la conversione della maggior parte degli algoritmi in OpenCL!
• Il tester non seleziona la CPU per OpenCL - questo non è riportato da nessuna parte!

Probabilmente sono il primo a voler accelerare il test a spese della GPU, avendo letto le promesse...

MetaDriver: Il primo thread su OpenCL descrive anche abbastanza bene i criteri che un algoritmo deve possedere per avere un potenziale di accelerazione ocl.

Rileggi il mio post.

Il criterio principale: l'esecuzione del codice MQL in "stile OpenCL" dovrebbe superare il tempo = Numero di_Buffer * 0,35300 ms in 1 tick.

Per scoprire la velocità dell'algoritmo in MQL con una precisione di microsecondi (1000 microsecondi = 1 millisecondo), si dovrà eseguire nel tester diverse volte e Total_Time / Number_of_Ticks (il mio post superiore).

Se non fosse per il ritardo del buffer, il mio codice passerebbe il test in ~30 secondi - che è ~2 volte più veloce del MQL "stile OpenCL" (55 secondi) e ~11 volte più veloce del codice normale (320 secondi).

Quali altri criteri ci sono?

[Sceptic Philozoff]

Roffild: L'affermazione è che non ci sono informazioni negli articoli:

• Latenza di lettura/scrittura anche di un piccolo buffer = 0,35300 ms - questo è ciò che rende la maggior parte degli algoritmi in OpenCL senza senso!

A giudicare dalla vostra esperienza con OpenCL, avrete già capito che non tutti gli algoritmi sono direttamente accelerati. Uno dei problemi principali qui è minimizzare gli accessi alla memoria globale.

A proposito, ora devo risolvere un problema simile con l'accesso casuale alla memoria globale del dispositivo (troppo privato questo accesso casuale, ed è un cazzo di overhead). Lo risolverò non appena avrò rimesso in sesto il mio cervello.

Il tester non seleziona la CPU per OpenCL - questo non è riportato da nessuna parte!

Scrivete al Service Desk e giustificate la necessità di tale caratteristica.

Se il tester non viene utilizzato, è già fatto (questa è la mia applicazione). E non ho ancora controllato il tester.

[Roffild]

Mathemat:

È già stato scritto che non tutti gli algoritmi sono direttamente accelerati. Devi usare il tuo cervello qui, e uno dei problemi principali è quello di ridurre al minimo gli accessi alla memoria globale.

Bene, ora devo risolvere un problema simile con l'accesso casuale alla memoria globale (questo accesso casuale è troppo frequente). Lo risolverò non appena avrò il cervello in funzione.

È il momento di usare il cervello perché 0,35300 ms si riferisce esattamente a clEnqueue[Read/Write]Buffer() e non agli accessi globali alla memoria all'interno del kernel.

Il secondo può essere risolto ottimizzando il kernel stesso, mentre il primo è un limite di ferro.