OpenCL：真のチャレンジ - ページ 5

[Ilyas]

SDへのリクエストを 作成し、コードを添付してください（PMを介して行うことができます）、私はあなたのコードを分析します。

[Roffild]

私のコードのどの部分に興味があるのでしょうか？異なるファイルに依存することが多いのですが。

今抱えている問題は、テスターの1tickでバッファの書き込みと読み出しを行うだけで、チェックには十分なのです。

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

スクリプトによる実行。

2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) Device #1 = AMD Phenom(tm) II X4 925 Processor
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) デバイス#0 = Cypress
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) OpenCL: GPUデバイス 'Cypress' が選択されました。
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) build log =.
2013.10.30 18:55:40 OpenCL_buffer_test (EURUSD,H1) バッファサイズ(バイト) =240

テスターで2013.01.09から2013.10.10まで、M5で「OHLC on M1」を使ってエキスパートを実行中です。

2013.10.30 19:01:44 Core 1 EURUSD,M5: 2013.01.09 00:00 から 2013.10.10 00:00 まで expertsOpenCL_buffer_test.ex5 のテストが開始されました。
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 Device #0 = Cypress
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 OpenCL: GPU デバイス 'Cypress' が選択されました。
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 ビルドログ＝。
2013.10.30 19:01:44 Core 1 2013.01.09 00:00:00 バッファサイズ（バイト） =240
2013.10.30 19:04:55 Core 1 EURUSD,M5: 1108637 ticks (55953 bars) generated within 192443 ms (total bars in history 131439, total time 192521 ms)
2013.10.30 19:04:55 Core 1 294 Mb メモリ使用量

テスターには1つのデバイスしかないことに注意してください。

もし

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

では

2013.10.30 19:16:00 Core 1 EURUSD,M5: 88218 ms以内に1108637ティック（55953バー）が発生（履歴の総バー数131439、総時間88297 ms）しました。

[Ilyas]

CLExecuteは ジョブの実行をキューに入れるだけで、終了を待たずにMQLプログラムに制御を戻すので、実際にはカーネルの終了を待つ関数である

[Ilyas]

具体的には、この例では、OpenCLを使う 利点が、バッファコピーのオーバーヘッドで食われてしまっている。

OHLCのデータで計算を行う必要がある場合は、あらかじめ大きめのバッファを作っておき、新しいデータが来たときだけ、新しい始まりとバッファの大きさをカーネルに伝えて上書きする、いわゆるスペアライティングが必須となるのです。

[Roffild]

OHLC転送を最適化できたとしても（CLSetKernelArgを 使用して最後のバーを転送します）、結果バッファを読み込む際にクラッシュします。

2013.10.31 19:24:13 Core 1 EURUSD,M5: 1108637 ticks (55953 bars) generated within 114489 ms (total bars in history 131439, total time 114598 ms)

(CLBufferWrite(hbuffer[0], price); のある行をIFの下に移動しました。）

[Sceptic Philozoff]

Roffild: え...GPUのOpenCLを使った高速化に関する記事は、現実的な課題に対処していないため、おとぎ話であることが判明しました :(

では、誰がそれを阻んでいるのでしょうか？おとぎ話にならないようなリアルなものを書いてきてください。しかし、加速が起こるような例を探してみてください。ここが一番難しいところです。

私の記事についてなら...まあ、入門書を書いていたんですけどね。そして、行列の掛け算はなかなか便利な操作です。

追伸：ちなみに、CPUがIntelの場合、その CPUでのx86コアのエミュレーションは他社製CPUよりもはるかに高速になります。それは、コアごとに計算し直した場合です。

HD5850：基本的にはかなりまともなカードですが、最近のカードはより良くなっています - フライが増えただけでなく、OpenCLの最適化により、です。例えば、グローバルメモリアクセスタイムを大幅に短縮することができます。

P.P.S. OpenCL は万能ではなく、ごく稀に著しく高速化することができる実行可能なツールです。また、それ以外のあまり便利でないケースでは、加速度があまり印象的でない--あったとしても。

[Vladimir Gomonov]

Roffild:
え...GPUでOpenCLを使った高速化に関する記事は、実際のタスクを扱っていないため、おとぎ話であることが判明しました :(

そうではありません。 おとぎ話では、「どんなアルゴリズムでもOpenCLで高速化できる」とされています。 どんなアルゴリズムでもというわけではありません。

OpenCLの最初のスレッドでは、アルゴリズムがOCLアクセラレーションの可能性を持つために持っていなければならない基準について、非常によく説明されています。

頑張ってください。

[Roffild]

計算速度ではなく、2倍の高速化（0.02ms vs 0.05ms）を実現しました。

記事には情報がない、という主張です。

• 小さなバッファでも読み書きのレイテンシは0.35300ms - これが、ほとんどのアルゴリズムのOpenCLへの変換を無効にしているのです
• テスターがOpenCLのCPUを選択しない - これはどこにも報告されていません!

お約束を読んで、GPUを犠牲にしてでもテストを早くしたいと思ったのは、たぶん私が初めてだと思います...。

MetaDriver: OpenCLの最初のスレッドでは、アルゴリズムがOCL加速の可能性を持つために持っていなければならない基準について、非常によく説明されています。

私の投稿をもう一度読んでみてください。

主な基準：「OpenCLスタイル」でのMQLコードの実行は、1tickで時間 =Number of_Buffers * 0.35300 msを 超える必要があります。

マイクロ秒（1000マイクロ秒＝1ミリ秒）の精度でMQLのアルゴリズムの速度を知るには、テスターとTotal_Time / Number_of_Ticks（私のトップポスト）で何回か実行する必要があります。

これは、「OpenCLスタイル」のMQL（55秒）の2倍、通常のコード（320秒）の11倍の速さです。

他にどんな基準があるのでしょうか？

[Sceptic Philozoff]

Roffild: 記事には情報がない、という主張です。

• 小さなバッファでも読み込み/書き込みのレイテンシは0.35300ms、これがOpenCLのほとんどのアルゴリズムを無意味にしているのです

OpenCLを扱った経験から判断すると、すべてのアルゴリズムが直接的に高速化されるわけではないことは、すでに理解されているはずです。ここで大きな問題となるのが、グローバルメモリアクセスの最小化です。

ちなみに、今はグローバルデバイスメモリへのランダムアクセスについても同様の問題を解決しなければならない（このランダムアクセスがプライベートすぎて、クソオーバーヘッドになっている）。頭の回転が戻り次第、解決していきます。

テスターがOpenCL用のCPUを選択しない - これはどこにも報告されていません!

サービスデスクに、そのような機能が必要であることを説明する。

テスターを使わなければ、もう終わりです（これは私の用途です）。そして、テスターでの確認はまだしていません。

[Roffild]

Mathemat:

すべてのアルゴリズムが直接的に高速化されるわけではないことは、すでに述べたとおりである。ここで頭を使わなければならないのですが、大きな問題のひとつは、グローバルメモリアクセスを最小限にすることです。

さて、今度はグローバルメモリへのランダムアクセス（このランダムアクセスが頻繁すぎる）についても、同様の問題を解決しなければなりません。頭を働かせたらすぐに解決します。

0.35300msはまさにclEnqueue[Read/Write]Buffer()のことで、カーネル 内のグローバルメモリアクセスのことではないので、頭を使うときが来たのでしょう。

2つ目はカーネル自体の最適化で解決できますが、1つ目は鉄の限界です。