OpenCL : de vrais défis - page 8

[Roffild]

Mathemat:

Regardez votre propre code : Et puis, à la dernière ligne, vous divisez vous-même 240 par 18 (ce sont des unités pour votre carte).

Vous êtes manifestement confus à propos de quelque chose. Voici l'article controversé :

   double price[30];

   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); /// <-- здесь же НЕ sizeof(price) вообще-то!
   local_work_size[0] = global_work_size[0] / units;

Print("global=",global_work_size[0]," local=",local_work_size[0]);

Conclusion : global=30 local=1

Et 240 octets, c'est exactement le moment où vous créez le tampon.

[Sceptic Philozoff]

Roffild:

Vous êtes manifestement confus à propos de quelque chose. Voici une pièce controversée :

Sortie : global=30 local=1

Et 240 octets exactement lors de la création du tampon.

global = 240. Imprimer

global_work_size[0]

Et local_work_size[0] = (uint) 240/18 = 13

P.S. Oui, vous avez raison. Pardon. J'ai été un peu confus.

local_work_size[0] = (uint) 30/18 = 1. Et j'ai la même chose, puisque les unités = 28.

[Sceptic Philozoff]

Encore une fois, Roffild:

Mathemat: Давай тупо прикинем. 18 задач, выполняемых одновременно на мухах GPU, - это максимум то, что можно сделать на 4-5 нитках CPU. А CPU на x86 эмуляции может организовать гораздо больше ниток. Во всяком случае, если это Intel. Мой бывший Pentium G840 (2 ядра) дал ускорение примерно в 70 раз - на двух unit'ах! Я уже не говорю о том, что вытворяет мой текущий... условно говоря, i7.

Une tâche bien parallélisée (voir les scripts de MetaDriver à partir du premier thread ocl) peut atteindre des accélérations jusqu'à 1000 ou plus sur GPU (par rapport à l'exécution d'un seul thread sur CPU sur MQL5). Si vous ne le trouvez pas - je peux vous l'envoyer, vous pouvez le tester sur votre carte.

[Roffild]

Avez-vous trouvé le tampon et sa vitesse ?

Vous feriez mieux d'utiliser AMD CodeXL pour déterminer les UNITÉS etc. - il a de beaux graphiques de performance.

AMD CodeXL lui-même est défaillant mais il est difficile de tirer des conclusions sans lui.

Je n'utiliserai pas OpenCL tant que le testeur ne me permettra pas d'utiliser le CPU ou tant que je n'exécuterai pas une tâche dont la durée est supérieure à Number of_buffers * 0,353 msec.

P.S.

J'ai finalement fini d'optimiser mon code et la variante finale passe le test en 33 secondes (320 secondes - avant optimisation, 55 secondes - "style OpenCL").

[Sceptic Philozoff]

Roffild: Avez-vous trouvé le tampon et sa vitesse ?

Il n'y a rien à découvrir. Il est clair qu'il s'agit d'une opération lente. Conclusion - augmenter le travail à l'intérieur du noyau (il y en a trop peu dans votre code).

Et achetez une carte vidéo plus moderne, il semble s'être amélioré avec elle.

AMD CodeXL lui-même est défaillant mais il est difficile de tirer des conclusions sans lui.

L'utilitaire d'Intel est plutôt utile aussi, mais pour les pierres Intel. Eh bien, et pour attraper les erreurs les plus évidentes dans le noyau.

P.S. J'ai, après tout, fini d'optimiser mon code et la variante finale passe maintenant le test en 33 secondes (320 secondes - avant optimisation, 55 secondes - "OpenCL-style").

C'est déjà beaucoup mieux.

[Roffild]

Aujourd'hui, j'ai eu besoin de générer un tableau avec 1 bit dans les chiffres.

En même temps, je me suis entraîné avec OpenCL.

Je poste le code pour démontrer une méthode intéressante de calcul de global_work_size et local_work_size. L'idée elle-même est tirée de IntrotoOpenCL.pdf (j'en ai une copie), mais je l'ai modifiée.

void OnStart()
{
   const string source =
      "kernel void differ(const int sizearray, const int bits, global uchar *result)        \r\n"
      "{                                                                                    \r\n"
      "   size_t global_index = get_global_id(0);                                           \r\n"
      "   if (global_index >= sizearray) return; // проверка границ, когда work>arraysize   \r\n"
      "   size_t xor = global_index;                                                        \r\n"
      "   uchar secc = 0;                                                                   \r\n"
      "   for (int bit = bits; bit>-1; bit--)                                               \r\n"
      "     if ((xor & ((size_t)1 << bit)) > 0) secc++;                                     \r\n"
      "   result[global_index] = secc;                                                      \r\n"
      "}                                                                                    \r\n"
   ;
   
   int hContext = CLContextCreate();
   string build_log = "";
   int hProgram = CLProgramCreate(hContext, source, build_log);
   Print("Error = ",build_log);
   int hKernel = CLKernelCreate(hProgram, "differ");
   
   uchar alldiff[1 << 17] = {0};
   CLSetKernelArg(hKernel, 0, ArraySize(alldiff));
   CLSetKernelArg(hKernel, 1, 17 /*bits*/);
   int hBuffer = CLBufferCreate(hContext, sizeof(alldiff), CL_MEM_WRITE_ONLY);
   CLSetKernelArgMem(hKernel, 2, hBuffer);
   CLBufferWrite(hBuffer, alldiff);
   
   /*uchar group_size[1024] = {0};
   uint deviceinfo_size = 0;
   CLGetDeviceInfo(hContext, CL_DEVICE_MAX_WORK_GROUP_SIZE, group_size, deviceinfo_size);
      for (int x = deviceinfo_size; x>=0; x--) Print(group_size[x]);
      Print("ch ",CharArrayToString(group_size));
   */ ///// CLGetDeviceInfo возвращает массив битов (шо за бред?)
   uint group_size = 256;
   
   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(alldiff);
   local_work_size[0] = global_work_size[0] / units;
   if (local_work_size[0] < 1) local_work_size[0] = 1;
   if (local_work_size[0] > group_size) local_work_size[0] = group_size;
   if (global_work_size[0] % local_work_size[0] != 0)
   {
      // увеличиваем global, чтобы global % local == 0
      // в самом kernel проверяется выход за границы
      global_work_size[0] = (int(global_work_size[0] / local_work_size[0]) +1) * local_work_size[0];
      // объяснение в
      // http://wiki.rac.manchester.ac.uk/community/OpenCL?action=AttachFile&amp;do=get&target=IntrotoOpenCL.pdf
   }
      Print("work=", global_work_size[0], " local=", local_work_size[0], " group=", group_size);
   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, alldiff);
   
   int hDump = FileOpen("alldiff.diff", FILE_ANSI|FILE_WRITE);
   for (int x = 0, xcount = ArraySize(alldiff); x < xcount; x++)
      FileWriteString(hDump, (string)alldiff[x]+",");
   FileClose(hDump);
   
   CLBufferFree(hBuffer);
   CLKernelFree(hKernel);
   CLProgramFree(hProgram);
   CLContextFree(hContext);
}