OpenCL : de vrais défis - page 5

[Ilyas]

Créez une demande au SD et joignez le code (vous pouvez le faire par MP), je vais analyser votre code.

[Roffild]

Quelle partie de mon code vous intéresse exactement ? J'ai beaucoup de dépendances sur différents fichiers.

Le problème que j'ai maintenant est que l'écriture et la lecture du tampon ne se fait qu'en 1 tick du testeur, et pour la vérification c'est suffisant :

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

Exécution par script :

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositif #1 = Processeur AMD Phenom(tm) II X4 925
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Dispositif n°0 = Cypress
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) OpenCL : périphérique GPU 'Cypress' sélectionné
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Taille du tampon en octets =240

Exécution de l'expert dans le testeur du 2013.01.09 au 2013.10.10 sur M5 avec "OHLC sur M1" :

2013.10.30 19:01:44 Core 1 EURUSD,M5 : test des experts\OpenCL_buffer_test.ex5 du 2013.01.09 00:00 au 2013.10.10 00:00 commencé
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 Dispositif #0 = Cypress
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 OpenCL : périphérique GPU 'Cypress' sélectionné
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 taille du tampon en octets =240
2013.10.30 19:04:55 Core 1 EURUSD,M5 : 1108637 ticks (55953 barres) générés en 192443 ms (total barres dans l'historique 131439, temps total 192521 ms)
2013.10.30 19:04:55 Core 1 294 Mb de mémoire utilisée

Notez qu'il n'y a qu'un seul appareil dans le testeur.

Si

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

puis

2013.10.30 19:16:00 Core 1 EURUSD,M5 : 1108637 ticks (55953 barres) générés en 88218 ms (total barres dans l'historique 131439, temps total 88297 ms)

[Ilyas]

Read est nécessaire ; en fait, c'est la fonction qui attend que le noyau se termine, puisque CLExecute ne fait que mettre le travail en file d'attente pour exécution, puis renvoie le contrôle au programme MQL sans attendre qu'il se termine.

[Ilyas]

Dans cet exemple, l'avantage de l'utilisation d'OpenCL est annihilé par la surcharge liée à la copie de la mémoire tampon.

S'il est nécessaire d'effectuer des calculs sur des données OHLC, il est alors impératif de procéder à une écriture parcimonieuse, c'est-à-dire de créer à l'avance un tampon plus grand et de n'écraser ces nouvelles données que lorsque celles-ci arrivent, en indiquant au noyau le nouveau début et la taille du tampon.

[Roffild]

Même si nous parvenons à optimiser le transfert OHLC (nous utiliserons CLSetKernelArg pour transférer la dernière barre), nous nous planterons toujours lors de la lecture du tampon de résultats :

2013.10.31 19:24:13 Core 1 EURUSD,M5 : 1108637 ticks (55953 barres) générés en 114489 ms (total barres dans l'historique 131439, temps total 114598 ms)

(déplacement de la ligne avec CLBufferWrite(hbuffer[0], prix) ; sous IF)

[Sceptic Philozoff]

Roffild: Eh... Les articles sur l'augmentation de la vitesse en utilisant OpenCL sur les GPU se sont avérés être un conte de fées car ils n'ont pas vraiment abordé les tâches à accomplir :(

Eh bien, qui vous empêche de le faire ? Va et écris quelque chose de réel qui ne soit pas un conte de fées. Mais essayez de trouver un exemple pour que l'accélération se produise. C'est la partie la plus difficile.

Si vous parlez de mes articles... Eh bien, j'écrivais un abécédaire. Et la multiplication de la matrice est une opération très utile.

P.S. À propos, si votre CPU est de type Intel, l'émulation des cœurs x86 sur celui-ci est beaucoup plus rapide que sur un CPU concurrent. C'est-à-dire si vous le recalculez par cœur.

HD5850 : à la base, une carte assez décente, mais les cartes modernes sont meilleures - non seulement en raison du plus grand nombre de mouches, mais aussi en raison des optimisations OpenCL. Par exemple, le temps d'accès à la mémoire globale est considérablement réduit.

P.P.S. OpenCL n'est pas une panacée, c'est un outil viable qui peut accélérer de manière significative dans certains cas rares. Et dans d'autres cas moins pratiques, l'accélération n'est pas si impressionnante - s'il y en a une.

[Vladimir Gomonov]

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

Pas du tout. Le conte de fées est que "tout algorithme peut être accéléré dans OpenCL". Pas n'importe quel algorithme.

Le premier fil de discussion sur OpenCL décrit même assez bien les critères qu'un algorithme doit posséder pour avoir un potentiel d'accélération ocl.

Bonne chance avec ça.

[Roffild]

La revendication ne concerne pas la vitesse de calcul - il y a un gain de vitesse de 2x (0,02 ms contre 0,05 ms).

L'argument est qu'il n'y a aucune information dans les articles :

• Latence de lecture/écriture de même un petit tampon = 0,35300 ms - c'est ce qui invalide la conversion de la plupart des algorithmes à OpenCL !
• Le testeur ne sélectionne pas le CPU pour OpenCL - cela n'est signalé nulle part !

Je suis probablement le premier à vouloir accélérer le test au détriment du GPU, après avoir lu les promesses...

MetaDriver: Le premier fil de discussion sur OpenCL décrit même assez bien les critères qu'un algorithme doit posséder pour avoir un potentiel d'accélération ocl.

Relisez mon message.

Le critère principal : l'exécution du code MQL dans le "style OpenCL" doit dépasser le temps = Number of_Buffers * 0.35300 ms en 1 tick.

Pour connaître la vitesse de l'algorithme dans MQL avec une précision de l'ordre de la microseconde (1000 microsecondes = 1 milliseconde), il faudra exécuter dans le testeur plusieurs fois et Total_Time / Number_of_Ticks (mon post du haut).

Sans le retard de la mémoire tampon, mon code passerait le test en ~30 secondes - c'est ~2 fois plus rapide que le MQL "style OpenCL" (55 secondes) et ~11 fois plus rapide que le code normal (320 secondes).

Quels sont les autres critères ?

[Sceptic Philozoff]

Roffild: L'argument est qu'il n'y a aucune information dans les articles :

• Latence de lecture/écriture d'un tampon, même petit, = 0,35300 ms - c'est ce qui rend la plupart des algorithmes d'OpenCL insensés !

À en juger par votre expérience en matière d'OpenCL, vous devez déjà avoir compris que tous les algorithmes ne sont pas directement accélérés. L'un des principaux problèmes ici est de minimiser les accès à la mémoire globale.

Au fait, je dois maintenant résoudre un problème similaire avec l'accès aléatoire à la mémoire globale du dispositif (trop privé cet accès aléatoire, et c'est une putain de surcharge). Je le résoudrai dès que j'aurai remis mon cerveau sur les rails.

Le testeur ne sélectionne pas le CPU pour OpenCL - cela n'est signalé nulle part !

Écrivez au Service Desk et justifiez le besoin d'une telle fonctionnalité.

Si le testeur n'est pas utilisé, c'est qu'il est déjà fait (c'est mon application). Et je n'ai pas encore vérifié le testeur.

[Roffild]

Mathemat:

Il a déjà été écrit que tous les algorithmes ne sont pas directement accélérés. Vous devez utiliser votre cerveau ici, et l'un des principaux problèmes est de minimiser les accès à la mémoire globale.

Eh bien, maintenant je dois résoudre un problème similaire avec l'accès aléatoire à la mémoire globale (cet accès aléatoire est trop fréquent). Je le résoudrai dès que mon cerveau fonctionnera.

Il est temps d'utiliser votre cerveau car 0,35300 ms se réfère exactement à clEnqueue[Read/Write]Buffer() et non aux accès globaux à la mémoire à l'intérieur du noyau.

La seconde peut être résolue en optimisant le noyau lui-même alors que la première est une limite de fer.