OpenCL : de vrais défis - page 4

[Vladimir Kazakov]

TheXpert:

La méthode la plus difficile est FeedPatterns.

À première vue, il n'y a aucun moyen de l'accélérer : il y a très peu de travail pour les noyaux par rapport à la quantité de données (beaucoup de données, peu de travail), tous les gains seront mangés par la copie dans les deux sens.

Vous pouvez essayer de demander explicitement au compilateur de le paralléliser dans VS ou de créer des threads parallèles pour le processeur.

[Vladimir Gomonov]

TheXpert:
Finalement, je vais le transférer sur MQL.

Est-ce strictement nécessaire ? Je voulais me pencher sur l'AMP depuis longtemps, et maintenant il y a une chance...

[Mykola Demko]

kazakov.v:

À première vue, il n'y a aucun moyen de l'accélérer : il y a très peu de travail pour les noyaux par rapport à la quantité de données (beaucoup de données, peu de travail), tout le gain sera mangé par la copie dans les deux sens.

Vous pouvez essayer d'indiquer explicitement au compilateur de le paralléliser dans VS ou de créer des threads parallèles pour le processeur lui-même.

Pourquoi, ces tâches sont parfaitement adaptées à 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;
  }

Et le truc vert peut le faire aussi.

[Mykola Demko]

Urain:

Pourquoi, ces tâches sont excellentes sur OpenCL

et tu peux faire les trucs verts aussi.

Nous devons faire un test d'implémentation et de comparaison en OpenCL et C++, et s'il y a une augmentation, alors tout traduire.

[Roffild]

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* indique-t-il la taille maximale du vecteur ou la taille optimale ?

Lorsque CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, double3 et double4 seront-ils déjà lents ?

[Sceptic Philozoff]

Roffild:

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* indique-t-il la taille maximale du vecteur ou la taille optimale ?

Lorsque CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2, double3 et double4 seront-ils déjà lents ?

1. Au maximum.

2- Le ralentissement ne sera probablement pas grave, mais il n'augmentera pas la vitesse d'exécution.

[Roffild]

MQL OpenCL est juste une enveloppe sur l'API originale.
J'ai besoin de réponses et de clarifications sur la mise en œuvre de ce wrapper.

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
S'agit-il en fait d'une file d'attente pour un dispositif, et non d'un contexte ?

Les tampons de lecture/écriture sont-ils synchrones ou asynchrones ?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - ici CL_TRUE ou CL_FALSE ?

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

CLExecute() rend-il le contrôle immédiatement ? Il ne le bloque pas le temps de son exécution ?
Il semble qu'il faille 2 à 40 ms pour le placer dans la file d'attente...

Voici maintenant la question principale :
Quand et dans quelles conditions clFinish() est-il appelé ? Et en raison de l'absence de clFinish(), il est difficile de faire la queue.

Et l'aide MQL ne décrit pas du tout CL_MEM_*_HOST_PTR, mais ils y sont présents.

J'ai finalement converti complètement mon indicateur au style OpenCL.
Exécution du test du 2013.01.09 au 2013.10.10 sur M5 avec "OHLC sur M1" :
320 secondes - avant traduction
55 secondes - Émulation de style OpenCL sur MQL5 :

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

Mais la course GPU a été frustrante pour moi :(
J'avais espéré exécuter le test en moins de 30 secondes mais j'ai reçu un retard total pour CLBufferWrite !

Chargement de la carte vidéo à 32% et passage du test en 1710 secondes sans CL_MEM_*_HOST_PTR
Chargement de la carte vidéo à 22% et réalisation d'un test en 740 secondes avec CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR et CL_MEM_USE_HOST_PTR résultent en CLExecute : 5109 (ERR_OPENCL_EXECUTE)

Alors comment échanger correctement les données ?

Et toujours aucun CPU n'est sélectionné pour les calculs dans le testeur.

Adaptateur vidéo = ATI Radeon HD 5850
Processeur = Processeur AMD Phenom(tm) II X4 925

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X)) ?
S'agit-il en fait d'une file d'attente pour un dispositif, et non d'un contexte ?

Oui, le contexte et la file d'attente sont créés par dispositif (des recherches ont montré qu'opencl ne fonctionne pas correctement avec plusieurs dispositifs différents).

Les tampons de lecture/écriture sont-ils synchrones ou asynchrones ?
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - ici CL_TRUE ou CL_FALSE ?

La lecture et l'écriture sont synchrones.

bool CLExecute(int kernel) = clEnqueueTask() ;
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel() ;
clEnqueueNativeKernel() - non implémenté.
CLExecute() rend-il le contrôle immédiatement ? Il ne le bloque pas le temps de son exécution ?

Oui

Et maintenant, la question principale :
Quand et dans quelles conditions clFinish() est-il appelé ? Et en raison de l'absence de clFinish(), il est difficile de former une file d'attente.

Non appelé, la lecture depuis la mémoire doit être utilisée.

Et il n'y a aucune description de CL_MEM_*_HOST_PTR dans l'aide MQL.

J'ai finalement converti complètement mon indicateur au style OpenCL.
Exécution du test du 2013.01.09 au 2013.10.10 sur M5 avec "OHLC sur M1" :
320 secondes - avant traduction
55 secondes - Émulation de style OpenCL sur MQL5 :
Mais la course GPU a été frustrante pour moi :(
J'espérais que le test se déroulerait en moins de 30 ms et j'ai obtenu un retard total pour CLBufferWrite !

Chargement de la carte vidéo à 32% et passage du test en 1710 secondes sans CL_MEM_*_HOST_PTR
Chargement de la carte vidéo à 22% et test en 740 secondes avec CL_MEM_ALLOC_HOST_PTR
CL_MEM_COPY_HOST_PTR et CL_MEM_USE_HOST_PTR résultent en CLExecute : 5109 (ERR_OPENCL_EXECUTE)

Alors comment échanger correctement les données ?

Les drapeaux CL_MEM_COPY_HOST_PTR et CL_MEM_USE_HOST_PTR ne sont pas actuellement supportés par le terminal (nous étudions ce problème).

Et toujours aucun CPU n'est sélectionné pour les calculs dans le testeur.

Avez-vous essayé de spécifier explicitement le périphérique CPU ?

[Roffild]

Pourquoi ne pas essayer de nous donner des tampons asynchrones et clFinish() ?

On suppose que c'est l'écriture synchrone qui ralentit, ce que même AMD CodeXL laisse entendre :

"clEnqueueWriteBuffer : Synchronisation inutile. Bloquer l'écriture"

Et dans le testeur de CPU, il n'est même pas possible de le sélectionner par numéro. Bogue n° 865549.

[Roffild]

Eh... Les articles sur l'augmentation de la vitesse en utilisant OpenCL sur GPU se sont avérés être un conte de fées car ils ne traitaient pas de tâches réelles :(

Ce mois-ci, j'ai écrit des milliers de lignes de code pour conquérir OpenCL.

Ainsi, pour déboguer OpenCL, j'ai dû émuler des fonctions de MQL pour les exécuter à travers AMD CodeXL en C/C++.

Je vais répéter les résultats du test de 2013.01.09 à 2013.10.10 sur M5 avec "OHLC sur M1" :
320 secondes - avant traduction
55 secondes - Émulation de style OpenCL sur MQL5

Le "style OpenCL" consiste à réduire au maximum le nombre d'appels CopyHigh/CopyTime/CopyOpen/.... et en augmentant la quantité de code pour traiter les tableaux après l'appel de ces fonctions.

Et ces calculs sont ce qui manque aux beaux articles sur OpenCL :

Résultat du test sans OpenCL :
Core 1 EURUSD,M5 : 1108637 ticks (55953 barres) générés en 55427 ms (total des barres dans l'historique 131439, temps total 55520 ms)
55427 ms / 1108637 tick = 0.04999 ms/tick - 1 tick par CPU (l'exécution sur OpenCL ne devrait pas dépasser ce temps)

C'est ce que j'ai obtenu en exécutant mon propre code en C/C++ et en le faisant passer par AMD CodeXL :
0.02000 ms - 0.05000 ms - exécution de mon noyau sur GPU
0.35300 ms - un appel à clEnqueueWriteBuffer pour 168 octets à 500KB/s
0.35300 ms - un appel clEnqueueWriteBuffer pour 3.445 KBytes avec 9.500 MBytes/s (le temps de transfert moyen est le même)

168 Octets est :
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};

Et j'ai obtenu 3,445 KByte à cause d'une erreur de calcul de la taille du tableau 21*168, mais même cela n'a pas affecté le temps de transfert.

En fin de compte, même si je parviens à optimiser mon noyau à 0,02000 ms, ce qui est effectivement ~2 fois plus rapide que la passe MQL habituelle (0,04999 ms), tout se résume à la vitesse de lecture/écriture du GPU (0,35300 ms - ~7 fois plus lente que le calcul MQL !).

Le CPU n'est pas sélectionné dans mon testeur pour OpenCL, je ne peux donc pas utiliser 3 autres cœurs vides...

P.S.
55 secondes n'est pas encore la limite de l'optimisation dans MQL, c'est juste une émulation OpenCL quand il n'y a pas de support :)