OpenCL: sfide reali - pagina 4

[Vladimir Kazakov]

TheXpert:

Il metodo più difficile è FeedPatterns.

A prima vista, non c'è modo di velocizzarlo: c'è pochissimo lavoro per i kernel rispetto alla quantità di dati (molti dati, poco lavoro), tutti i guadagni saranno mangiati dalla copia avanti e indietro.

Potete provare a chiedere esplicitamente al compilatore di parallelizzarlo in VS o creare thread paralleli per il processore.

[Vladimir Gomonov]

TheXpert:
Alla fine, ho intenzione di trasferirlo su MQL.

È strettamente necessario? È da molto tempo che voglio dare un'occhiata a AMP, e ora c'è una possibilità...

[Mykola Demko]

kazakov.v:

A prima vista, non c'è modo di velocizzarlo: c'è pochissimo lavoro per i kernel rispetto alla quantità di dati (molti dati, poco lavoro), tutto il guadagno sarà mangiato dalla copia avanti e indietro.

Potete provare a dire esplicitamente al compilatore di parallelizzarlo in VS o creare thread paralleli per il processore stesso.

Perché, questi compiti sono perfettamente adatti 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;
  }

E anche la roba verde può farlo.

[Mykola Demko]

Urain:

Perché, questi compiti sono grandiosi su OpenCL

e puoi fare anche le cose verdi.

Bisogna fare un test di implementazione e confronto in OpenCL e C++, e se c'è un aumento, si può tradurre tutto.

[Roffild]

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* indica la dimensione massima del vettore o quella ottimale?

Quando CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, double3 e double4 saranno già lenti?

[Sceptic Philozoff]

Roffild:

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* indica la dimensione massima del vettore o quella ottimale?

Quando CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, double3 e double4 saranno già lenti?

1. Al massimo.

2- È improbabile che il rallentamento sia grave, ma non aumenterà la velocità di esecuzione.

[Roffild]

MQL OpenCL è solo un wrapper sopra l'API originale.
Ho bisogno di sapere risposte e chiarimenti sull'implementazione di questo wrapper.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Si tratta in realtà di una coda per un dispositivo e non di un contesto?

I buffer di lettura/scrittura sono sincroni o asincroni?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - qui CL_TRUE o CL_FALSE ?

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

CLExecute() restituisce immediatamente il controllo? Non lo blocca per il tempo della sua esecuzione?
Sembra che ci vogliano 2-40 ms per metterlo in coda...

Ora ecco la domanda principale:
Quando e in quali condizioni viene chiamata clFinish()? E a causa dell'assenza di clFinish(), è difficile da accodare.

E l'aiuto MQL non descrive affatto CL_MEM_*_HOST_PTR, ma sono presenti.

Ho finalmente convertito completamente il mio indicatore in stile OpenCL.
Esecuzione del test dal 2013.01.09 al 2013.10.10 su M5 con "OHLC su M1":
320 secondi - prima della traduzione
55 secondi - Emulazione in stile OpenCL su MQL5:

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

Ma la corsa della GPU è stata frustrante per me :(
Avevo sperato di eseguire il test in meno di 30 secondi ma ho ricevuto un ritardo totale per CLBufferWrite!

Caricamento della scheda video al 32% e superamento del test in 1710 secondi senza CL_MEM_*_HOST_PTR
Caricare la scheda video al 22% e fare un test in 740 secondi con CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR e CL_MEM_USE_HOST_PTR risultano in CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Quindi come scambiare correttamente i dati?

E ancora nessuna CPU è selezionata per i calcoli nel tester.

Adattatore video = ATI Radeon HD 5850
Processore = Processore AMD Phenom(tm) II X4 925

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
Si tratta in realtà di una coda per un dispositivo e non di un contesto?

Sì, il contesto e la coda sono creati per dispositivo (la ricerca ha dimostrato che opencl non funziona correttamente con diversi dispositivi).

I buffer di lettura/scrittura sono sincroni o asincroni?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - qui CL_TRUE o CL_FALSE ?

La lettura e la scrittura sono sincrone.

bool CLExecute(int kernel) = clEnqueueTask();
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel();
clEnqueueNativeKernel() - non implementato.
CLExecute() restituisce immediatamente il controllo? Non lo blocca per il tempo della sua esecuzione?

Sì

E ora la domanda principale:
Quando e in quali condizioni viene chiamata clFinish()? E a causa dell'assenza di clFinish(), è difficile formare una coda.

Non chiamato, deve essere usata la lettura dalla memoria.

E non ci sono descrizioni di CL_MEM_*_HOST_PTR nella Guida MQL.

Ho finalmente convertito completamente il mio indicatore in stile OpenCL.
Esecuzione del test dal 2013.01.09 al 2013.10.10 su M5 con "OHLC su M1":
320 secondi - prima della traduzione
55 secondi - Emulazione in stile OpenCL su MQL5:
Ma la corsa della GPU è stata frustrante per me :(
Speravo che il test venisse eseguito in meno di 30 ms e ho ottenuto un ritardo totale per CLBufferWrite!

Caricamento della scheda video al 32% e superamento del test in 1710 secondi senza CL_MEM_*_HOST_PTR
Caricamento della scheda video al 22% e test in 740 secondi con CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR e CL_MEM_USE_HOST_PTR risultano in CLExecute: 5109 (ERR_OPENCL_EXECUTE)

Quindi come scambiare correttamente i dati?

I flag CL_MEM_COPY_HOST_PTR e CL_MEM_USE_HOST_PTR non sono attualmente supportati dal terminale (stiamo indagando su questo problema).

E ancora nessuna CPU è selezionata per i calcoli nel tester.

Avete provato a specificare esplicitamente il dispositivo della CPU?

[Roffild]

Che ne dite di provare a darci buffer asincroni e clFinish()?

C'è l'ipotesi che sia la scrittura sincrona a rallentare, cosa a cui anche AMD CodeXL accenna:

"clEnqueueWriteBuffer: sincronizzazione non necessaria. Blocco della scrittura"

E nel tester della CPU non è nemmeno selezionabile per numero. Bug #865549.

[Roffild]

Eh... Gli articoli sull'aumento di velocità usando OpenCL su GPU si sono rivelati una favola perché non si sono occupati di compiti reali :(

Questo mese, ho scritto migliaia di righe di codice per conquistare OpenCL.

Così, per eseguire il debug di OpenCL, ho dovuto emulare le funzioni da MQL per eseguirle attraverso AMD CodeXL in C/C++.

Ripeterò i risultati del test dal 2013.01.09 al 2013.10.10 su M5 con "OHLC su M1":
320 secondi - prima della traduzione
55 secondi - emulazione in stile OpenCL su MQL5

Lo "stile OpenCL" consiste nel ridurre al minimo il numero di chiamate CopyHigh/CopyTime/CopyOpen/.... e aumentando la quantità di codice per elaborare gli array dopo che queste funzioni sono chiamate.

E questi calcoli sono ciò che manca ai bei articoli su OpenCL:

Risultato del test senza OpenCL:
Core 1 EURUSD,M5: 1108637 ticks (55953 barre) generati in 55427 ms (barre totali nella storia 131439, tempo totale 55520 ms)
55427 ms / 1108637 tick = 0.04999 ms/tick - 1 tick per CPU (l'esecuzione su OpenCL non dovrebbe superare questo tempo)

Questo è quello che ho ottenuto eseguendo il mio codice in C/C++ ed eseguendolo attraverso AMD CodeXL:
0.02000 ms - 0.05000 ms - esecuzione del mio kernel su GPU
0.35300 ms - una chiamata a clEnqueueWriteBuffer per 168 byte a 500KB/s
0.35300 ms - una chiamata clEnqueueWriteBuffer per 3.445 KBytes con 9.500 MBytes/s (il tempo medio di trasferimento è lo stesso)

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 ho ottenuto 3.445 KByte a causa di un errore di calcolo della dimensione dell'array 21*168, ma anche questo non ha influenzato il tempo di trasferimento.

In conclusione: anche se riesco a ottimizzare il mio kernel a 0,02000 ms, che è effettivamente ~2 volte più veloce del solito passaggio MQL (0,04999 ms), tutto si riduce alla velocità di lettura/scrittura della GPU (0,35300 ms - ~7 volte più lento del calcolo MQL!)

La CPU non è selezionata nel mio tester per OpenCL, quindi non posso usare altri 3 core vuoti...

P.S.
55 secondi non è ancora il limite dell'ottimizzazione in MQL, è solo un'emulazione OpenCL quando non c'è supporto :)