OpenCL：真正的挑战 - 页 8

[Roffild]

Mathemat:

看看你自己的代码。然后，在最后一行，你自己用240除以18（这是你的卡片的单位）。

你显然对某些事情感到困惑。以下是有争议的文章。

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

结论：global=30 local=1

而240字节正好是你创建缓冲区的时间。

[Sceptic Philozoff]

Roffild:

你显然对某些事情感到困惑。这里有一篇有争议的文章。

输出：global=30 local=1

而在创建缓冲区时正好是240字节。

全局=240。打印出来

global_work_size[0]

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

P.S. 是的，你说对了。请原谅。我有点糊涂了。

local_work_size[0] = (uint) 30/18 = 1.而我也一样，因为单位=28。

[Sceptic Philozoff]

还是那句话，罗菲德。

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

一个良好的并行化的任务（见MetaDriver 的第一个ocl线程的脚本）在GPU上可以实现高达1000以上的速度（与MQL5在CPU上的1线程执行相比）。如果你找不到它--我可以把它寄给你，你可以在你的卡上测试。

[Roffild]

你已经弄清楚了缓冲器和它的速度吗？

你最好使用AMD CodeXL来计算UNITS等--它有漂亮的性能图。

AMD CodeXL本身有故障，但没有它就很难得出任何结论。

在测试者允许我使用CPU之前，或者在我运行一个持续时间超过Number of_buffers * 0.353 msec的任务之前，我不会使用OpenCL。

P.S.

我确实完成了对代码的优化，最终的变体在33秒内通过了测试（320秒--优化前，55秒--"OpenCL式"）。

[Sceptic Philozoff]

Roffild: 你已经弄清楚了缓冲器和它的速度吗？

没有什么可琢磨的。很明显，这是一个缓慢的行动。结论--增加内核 内的工作（你的代码中的工作太少了）。

并买一块更现代的显卡，它似乎已经变得更好了。

AMD CodeXL本身有故障，但没有它就很难得出任何结论。

英特尔的工具也相当有用，但对英特尔的石头来说。嗯，还有就是为了捕捉内核中最明显的错误。

P.S. 我毕竟完成了对代码的优化，最终的变体在33秒内通过了测试（320秒--优化前，55秒--"OpenCL式"）。

已经好多了。

[Roffild]

今天我需要生成一个数字为1比特的数组。

同时，我还用OpenCL进行了练习。

我发布这段代码是为了演示一种有趣的计算 global_work_size 和 local_work_size 的方法。这个想法本身取自IntrotoOpenCL.pdf（我有一份副本），但我对它进行了调整。

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