OpenCL: gerçek problemler - sayfa 8

[Roffild]

Mathemat :

Pekala, kendi kodunuza bakın: Ve son satırda, 240'ı kendiniz 18'e bölersiniz (bunlar haritanız için birimlerdir).

Bir şeyi açıkça karıştırıyorsunuz. İşte tartışmalı kısım:

   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]);

Çıktı: global=30 yerel=1

Ve bir arabellek oluştururken 240 bayt.

[Sceptic Philozoff]

Roffild :

Bir şeyi açıkça karıştırıyorsunuz. İşte tartışmalı kısım:

Çıktı: global=30 yerel=1

Ve bir arabellek oluştururken 240 bayt.

genel = 240. Yazdır

global_work_size[ 0 ]

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

PS Evet, haklısınız. Afedersiniz. Biraz kafam karıştı.

local_work_size[0] = (uint) 30/18 = 1 birim=28.

[Sceptic Philozoff]

Bir kez daha, Roffield :

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

İyi paralel bir görev ( MetaDriver 'ocl ile ilgili ilk daldan bir komut dosyasına bakın), GPU'da 1000 veya daha fazla hızlanma elde etmenize olanak tanır (MQL5'te CPU üzerinde 1 iş parçacığında yürütmeye kıyasla). Bulamazsan atarım, haritanda test ederim.

[Roffild]

Tamponu ve hızını anladınız mı?

Ve UNITS, vb. hakkında AMD CodeXL ile çözmek daha iyi - performans grafikleri orada güzel.

AMD CodeXL'in kendisi buggy bir enfeksiyondur, ancak onsuz herhangi bir sonuç çıkarmak zordur.

Test cihazı CPU'ya izin verene veya arabellek sayısı * 0,353 ms'den daha uzun süren bir görev gelene kadar artık OpenCL ile uğraşmayacağım.

not

Sonunda kodumu optimize etmeyi bitirdim ve şimdi son sürüm testi 33 saniyede geçiyor (320 saniye - optimizasyondan önce, 55 saniye - "OpenCL tarzında").

[Sceptic Philozoff]

Roffild : ?

Ve bir şeyi anlamak için orada ne var. Bunun yavaş bir işlem olduğu açıktır. Sonuç, çekirdeğin içindeki işi arttırmaktır (kodunuzda çok az şey var).

Ve daha modern vidyahi alın, bununla daha iyi hale geldiler.

AMD CodeXL'in kendisi buggy bir enfeksiyondur, ancak onsuz herhangi bir sonuç çıkarmak zordur.

Intel yardımcı programı da çok kullanışlıdır - ancak Intel taşları için. Peki, çekirdekteki en bariz hataları yakalamak için.

PS Sonunda kodumu optimize etmeyi bitirdim ve şimdi son sürüm testi 33 saniyede geçiyor (320 saniye - optimizasyondan önce, 55 saniye - "OpenCL tarzında").

Zaten çok daha iyi.

[Roffild]

Bugün sayılarda 1 bitlik bir dizi oluşturmam gerekiyordu.

Aynı zamanda OpenCL üzerinde çalıştım.

Kodu global_work_size ve local_work_size hesaplamak için ilginç bir yöntemin gösterimi olarak gönderiyorum. Bu fikrin kendisi IntrotoOpenCL.pdf'den alınmıştır (bir kopyam var), ama düzelttim.

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