3

私はOpenCLを調べていますが、このカーネルが実行されると予想される方法と比較して、なぜこのカーネルの実行が非常に遅いのか少し混乱しています。カーネルは次のとおりです。

__kernel void copy(
  const __global char* pSrc, 
  __global __write_only char* pDst, 
  int length)
{
  const int tid = get_global_id(0);

  if(tid < length) {
    pDst[tid] = pSrc[tid];
  }
}

次の方法でバッファを作成しました。

char* out = new char[2048*2048];
cl::Buffer(
  context,
  CL_MEM_USE_HOST_PTR | CL_MEM_WRITE_ONLY,
  length,
  out);

入力バッファについても同様ですが、inポインタをランダムな値に初期化した点が異なります。最後に、カーネルを次のように実行します。

cl::Event event;
queue.enqueueNDRangeKernel(
  kernel, 
  cl::NullRange,
  cl::NDRange(length),
  cl::NDRange(1), 
  NULL, 
  &event);

event.wait();

次のように計算すると、平均時間は約75ミリ秒です。

cl_ulong startTime = event.getProfilingInfo<CL_PROFILING_COMMAND_START>();
cl_ulong endTime = event.getProfilingInfo<CL_PROFILING_COMMAND_END>();
std::cout << (endTime - startTime) * SECONDS_PER_NANO / SECONDS_PER_MILLI << "\n";

Intel i5-3450チップ（Sandy Bridgeアーキテクチャ）を搭載したWindows7を実行しています。比較のために、コピーを行う「直接」の方法は5ミリ秒未満かかります。event.getProfilingInfoには、ホストとデバイス間の通信時間が含まれているとは思いません。考え？

編集：

ananthonlineの提案で、charsの代わりにfloat4sを使用するようにカーネルを変更しました。これにより、平均実行時間が約50ミリに短縮されました。それでも私が望んでいたほど速くはありませんが、改善されています。ananthonlineに感謝します！

4

3 に答える 3

3

あなたの主な問題は、使用している 2048*2048 ワークグループだと思います。システム上の opencl ドライバーは、単一項目のワーク グループが多数ある場合、さらに多くのオーバーヘッドを管理する必要があります。GPU を使用してこのプログラムを実行すると、ハードウェアの飽和レベルが非常に低くなるため、これは特に悪いことです。

最適化: より大きなワーク グループでカーネルを呼び出します。既存のカーネルを変更する必要さえありません。質問を参照してください: このサイズはどのくらいですか? 例として以下の 64 を使用しました。ほとんどのハードウェアでは、たまたま 64 が妥当な数です。

cl::size_t myOptimalGroupSize = 64;
cl::Event event;
queue.enqueueNDRangeKernel(
  kernel, 
  cl::NullRange,
  cl::NDRange(length),
  cl::NDRange(myOptimalGroupSize), 
  NULL, 
  &event);

event.wait();

また、単一の値をコピーする以上のことをカーネルに行わせる必要があります。ここで、グローバルメモリに関する同様の質問に回答しました。

于 2012-10-11T22:57:33.980 に答える
1

CPU は GPU とは大きく異なります。これを x86 CPU で実行する場合、まともなパフォーマンスを達成するための最良の方法は、char または float4 の代わりに double16 (最大のデータ型) を使用することです (他の人が提案したように)。

CPU で OpenCL を使用した経験はほとんどありませんが、OpenMP の並列化で得られるパフォーマンス レベルに到達したことはありません。CPU と並行してコピーを行う最善の方法は、コピーするブロックを少数の大きなサブブロックに分割し、各スレッドにサブブロックをコピーさせることです。GPU のアプローチは直交的です。各スレッドは同じブロックのコピーに参加します。これは、GPU では、異なるスレッドが連続したメモリ領域に効率的にアクセスできる (合体) ためです。

OpenCL を使用して CPU で効率的なコピーを行うには、カーネル内でループを使用して連続したデータをコピーします。次に、使用可能なコア数以下のワークグループ サイズを使用します。

于 2012-10-12T08:31:15.710 に答える
1

cl::NDRange(1) がランタイムに単一アイテムのワーク グループを使用するように指示していると思います。これは効率的ではありません。C API では、これに NULL を渡して、ワーク グループのサイズをランタイムに任せることができます。C++ API にもそれを行う方法があるはずです (おそらく NULL だけでも)。これは CPU でより高速になるはずです。それは確かにGPU上にあります。

于 2012-10-14T23:30:57.877 に答える