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バイトを正確に。

global = 240です。プリントアウト

global_work_size[0]

そして、local_work_size[0] = (uint) 240/18 = 13となります。

追伸：はい、正解です。失礼しました。ちょっと混乱しました。

local_work_size[0] = (uint) 30/18 = 1.そして、units=28なので、私も同じです。

[Sceptic Philozoff]

再び、ロフィルド。

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

うまく並列化されたタスク（最初のoclスレッドからMetaDriverの スクリプトを参照）は、GPU上で1000以上のスピードアップを達成できます（MQL5でCPU上の1スレッド実行と比較）。もし見つからなかったら-私が送りますので、あなたのカードでテストしてみてください。

[Roffild]

バッファとその速度は把握されていますか？

UNITSなどを把握するには、AMD CodeXLを使ったほうがいいですね。性能グラフがきれいに表示されますよ。

AMD CodeXL自体に不具合がありますが、これがないと結論を出すのは難しいです。

テスターがCPUを使えるようにするか、Number of_buffers * 0.353 msecより長く続くタスクができるまで、OpenCLを使うつもりはないです。

追伸

私は最終的にコードの最適化を終え、最終的な変形は33秒でテストに合格しました（320秒-最適化前、55秒-「OpenCL-style」）。

[Sceptic Philozoff]

Roffild: バッファーとその速度は把握できましたか？

解明することは何もない。動作が遅いのは明らかです。結論 -カーネル 内の作業を増やす (あなたのコードにはそれが少なすぎる)。

そしてもっと最新のビデオカードを買ってください、それで良くなっているようです。

AMD CodeXL自体は不具合がありますが、これがないと結論を出すのは難しいです。

Intelのユーティリティも割と便利なんだけど、Intelの石のために。それと、カーネルの最も明白なエラーを検出するためです。

追伸：結局、私のコードは最適化され、最終的なバリエーションは33秒でテストに合格しました（320秒 - 最適化前、55秒 - 「OpenCL-style」）。

もう、だいぶ良くなりました。

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