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

[Vladimir Kazakov]

TheXpert:

一番難しいのはFeedPatternsという方法です。

一見すると、高速化する方法はありません。データ量に対してカーネルの 仕事は非常に少なく（多くのデータ、少しの仕事）、すべての利益は前後にコピーすることで食いつぶしてしまうでしょう。

VSでコンパイラに明示的に並列化を依頼したり、プロセッサに並列スレッドを作成してみたりするのも良いでしょう。

[Vladimir Gomonov]

TheXpert:
最終的には、MQLに移管するつもりです。

以前からAMPについて調べてみたいと思っていたのですが、今回チャンスがあったので......。

[Mykola Demko]

kazakov.v:

一見すると、高速化する方法はありません。データ量に対してカーネルの 仕事は非常に少なく（多くのデータ、少しの仕事）、すべての利得は前後にコピーすることで食べられてしまいます。

VSでコンパイラに明示的に並列化を指示したり、プロセッサに並列スレッドを作成してみたりするのも良いでしょう。

なぜかというと、これらのタスクはOpenCLに完全に適しているのです。

void Solver::FeedPattern(const Pattern &pattern)
  {
   size_t size=pattern.Size();

   assert(size==m_PatternSize);
   if(size!=m_PatternSize)
     {
      return;
     }

   const std::vector<double>&values=pattern.Values();
   double etalon=pattern.Etalon();

   size_t i;

   for(i=0; i<size;++i)
     {
      for(size_t j=0; j<size;++j)
        {
         m_Matrix[i][j]+=values[i]*values[j];
        }
      m_Matrix[i][size]+=values[i];
      m_Matrix[i][size+1]+=values[i]*etalon;
     }

   for(i=0; i<size;++i)
     {
      m_Matrix[size][i]+=values[i];
     }
   m_Matrix[size][size]+=1;
   m_Matrix[size][size+1]+=etalon;
  }

そして、緑色のものもできるのです。

[Mykola Demko]

Urain:

なぜかというと、これらのタスクはOpenCLでは素晴らしいことなのです

とか、緑色のものもできるんですね。

OpenCLとC++で実装と比較テストをして、増えるようなら全部翻訳する必要がありますね。

[Roffild]

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* は、最大ベクターサイズか最適なサイズかを示しますか？

CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 のとき、double3 と double4 はすでに遅いのでしょうか？

[Sceptic Philozoff]

Roffild:

CL_DEVICE_PREFERRED_VECTOR_WIDTH_* は、最大ベクターサイズか最適なサイズかを示しますか？

CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE=2 のとき、double3 と double4 はすでに遅いのでしょうか？

1.最大にする。

2- 遅くなることは少ないが、実行速度が上がらない。

[Roffild]

MQL OpenCLは、オリジナルのAPIのラッパーに過ぎません。
このラッパーの実装について、回答や説明を知りたいのですが。

CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X))?
これは実際にはコンテキストではなく、1つのデバイスのキューなのでしょうか？

バッファの読み書きが同期か非同期か？
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - ここでCL_TRUEまたはCL_FALSE?

bool CLExecute(intkernel) = clEnqueueTask()です。
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel()．
clEnqueueNativeKernel() - 未実装です。

CLExecute()は、すぐに制御を戻すのでしょうか？実行時間中はブロックされないのですか？
キューにセットするのに2～40msかかるようです...。

さて、ここからが本題です。
clFinish()は、いつ、どのような条件で呼び出されるのですか？ また、clFinish()がないため、キューイングが困難です。

また、MQLのヘルプにはCL_MEM_*_HOST_PTRの記述が全くありませんが、そこには存在します。

ようやく自分のインジケータをOpenCLスタイルに完全変換しました。
2013.01.09から2013.10.10まで、M5で「OHLC on M1」を使ってテストを実行中。
320秒 - 翻訳前
55秒 - MQL5でOpenCLスタイルのエミュレーションを行います。

// подготовка данных общая и копия kernel на MQL5, а эмуляция через:
for (int get_global_id = maxcount-1; get_global_id>-1; get_global_id--) NoCL(params,m_result,get_global_id);

でも、GPUの走りは、私にとっては悔しいものでした :(
30秒以内にテストを実行したかったのですが、CLBufferWriteでトータルラグを受け取りました

ビデオカードを32%でロードし、CL_MEM_*_HOST_PTRなしで1710秒でテスト通過
ビデオカードを22%でロードし、CL_MEM_ALLOC_HOST_PTRで740秒でテスト作成
CL_MEM_COPY_HOST_PTR と CL_MEM_USE_HOST_PTR は CLExecute: 5109 (ERR_OPENCL_EXECUTE) となる。

では、どのようにすれば適切にデータをやり取りできるのでしょうか。

そして、やはりテスターで計算するCPUは選択されません。

ビデオアダプタ＝ATI Radeon HD 5850
プロセッサー＝AMD Phenom(tm) II X4 925 プロセッサー

[Ilyas]

Roffild:
CLContextCreate() = clCreateCommandQueue(clCreateContextFromType(CL_DEVICE_TYPE_X))?
これは実際にはコンテキストではなく、1つのデバイスのキューなのでしょうか？

はい、コンテキストとキューはデバイスごとに作成されます（調査により、openclは複数の異なるデバイスで正しく動作しないことが判明しています）。

バッファの読み書きが同期か非同期か？
clEnqueue[Read/Write]Buffer(enqueue, buffer, CL_TRUE) - ここでCL_TRUEまたはCL_FALSE?

読み出しと書き込みは同期です。

bool CLExecute(intkernel) = clEnqueueTask()です。
bool CLExecute(int kernel, uint work_dim, works...) = clEnqueueNDRangeKernel()．
clEnqueueNativeKernel() - 未実装です。
CLExecute()は、すぐに制御を戻すのでしょうか？実行時間中はブロックされないのですか？

はい

そして、いよいよ本題です。
clFinish()は、いつ、どのような条件で呼び出されるのですか？ また、clFinish()がないため、キューを形成することが困難である。

呼び出されず、メモリからの読み込みを使用する必要があります。

また、MQLヘルプにはCL_MEM_*_HOST_PTRの記述が全くありません。

ようやく自分のインジケータをOpenCLスタイルに完全変換しました。
2013.01.09から2013.10.10まで、M5で「OHLC on M1」を使ってテストを実行中。
320秒 - 翻訳前
55秒 - MQL5でOpenCLスタイルのエミュレーションを行います。
でも、GPUの走りは、私にとっては悔しいものでした :(
30ms以下でテストが実行されることを期待していたのですが、CLBufferWriteのために完全なラグが発生してしまいました

ビデオカードを32%でロードし、CL_MEM_*_HOST_PTRなしで1710秒でテスト通過
ビデオカードを22%でロードし、CL_MEM_ALLOC_HOST_PTRを使用して740秒でテストを実施。
CL_MEM_COPY_HOST_PTR と CL_MEM_USE_HOST_PTR は CLExecute: 5109 (ERR_OPENCL_EXECUTE) の結果となる。

では、どのようにすれば適切にデータをやり取りできるのでしょうか。

CL_MEM_COPY_HOST_PTR と CL_MEM_USE_HOST_PTR フラグは、現在ターミナルでサポートされていません (この問題は調査中です)。

そして、やはりテスターで計算するCPUは選択されません。

CPUデバイスの明示的な指定は試されましたか？

[Roffild]

非同期バッファとclFinish()を提供してみるのはどうでしょう？

遅くなるのは同期書き込みであるという前提があり、それはAMD CodeXLでも示唆されている。

"clEnqueueWriteBuffer: Unnecessary synchronization.ブロッキングライト"

しかも、CPUテスターでは番号で選択することもできない。バグ番号: 865549

[Roffild]

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

今月は、OpenCLを攻略するために何千行もコードを書きました。

したがって、OpenCLをデバッグするためには、MQLから関数をエミュレートして、C/C++のAMD CodeXLで実行する必要があったのです。

2013.01.09～2013.10.10までのテスト結果をM5で「OHLC on M1」を使って繰り返してみます。
320秒 - 翻訳前
55秒 - MQL5でOpenCLスタイルのエミュレーションを行う。

"OpenCLスタイル "とは、CopyHigh/CopyTime/CopyOpen/...の呼び出し 回数を極力減らすことです。と、これらの関数が呼ばれた後に配列を処理するコード量を増やしています。

そして、この計算こそが、OpenCLの美談に欠けているものなのです。

OpenCLを使用しない場合のテスト結果。
Core 1 EURUSD,M5: 55427 ms以内に1108637ティック（55953バー）が発生（履歴の総バー数131439、総時間55520ms）。
55427 ms / 1108637 tick = 0.04999 ms/tick - 1CPUあたり1tick（OpenCLでの実行はこの時間を超えない ようにすること）

これは、C/C++で書いた自分のコードをAMD CodeXLで実行した結果です。
0.02000 ms - 0.05000 ms - GPU上での私のカーネルの実行
0.35300 ms - clEnqueueWriteBufferの1回の呼び出し（168バイト、500KB/sの場合
0.35300 ms - clEnqueueWriteBufferを1回呼び出し、3.445 KBytes、9.500 MBytes/s (平均転送時間は同じです)

168Byteです。
double open[21]={1.3668,1.3661,1.36628,1.3664,1.36638,1.36629,1.3664,1.36881,1.36814,1.3692,1.36918,1.36976,1.36816,1.36776,1.36779,1.3695,1.36927,1.36915,1.3679,1.36786,1.36838};

そして、21*168の配列サイズの計算ミスで3,445KByteになりましたが、それさえも転送時間には影響しませんでした。

結論： カーネルを0.02000msに最適化できたとしても、それは通常のMQLパス（0.04999ms）よりも確かに2倍速いのですが、すべてはGPUの読み取り/書き込み速度（0.35300ms - MQL計算の7倍遅い！）に帰結してしまいます。

テスターでOpenCLのCPUが選択されていないので、あと3コアの空きが使えない...。

追伸
55秒はまだMQLでの最適化の限界ではなく、サポートがないときのOpenCLのエミュレーションに過ぎません :)