0

これが私の問題です。プロジェクトを高速化するために、カーネル内で生成された値を共有メモリに保存したいのですが、その値を保存するのに非常に長い時間がかかることがわかりました。「THIS LINE」を削除すると（以下のコードを参照）、つまり「THIS LINE」を削除すると、その値を保存するのが非常に高速になります（100倍のスピードアップ！）。

extern __shared__ int sh_try[];

__global__ void xxxKernel (...)
{
  float v, e0, e1;
  float t;
  int count(0);
  for (...)
  {
     v = fetchTexture();
     e0 = fetchTexture();
     e1 = fetchTexture();
     t = someDeviceFunction(v, e0, e1);
     if (t>0.0 && t < 1.0)  <========== <THIS LINE>
       count++;
  }
  sh_try[threadIdx.x] = count;
}

main()
{
  sth..
  START TIMING:

  xxxKernel<<<gridDim.x, BlockDim.x, BlockDim.x*sizeof(int)>>> (...);

  cudaDeviceSynchronize();

  END TIMING.
  sth...
}

この問題を解決するために、データを共有メモリに保存するだけのコードを単純化します。停止します。私が知っているように、共有メモリ。最も効率的なメモリです。登録以外に、この高いレイテンシーは正常なのか、それとも私が間違ったことをしたのだろうかと思います。アドバイスをください!!! よろしくお願いします！

トルディ

更新: 共有メモリをグローバルメモリに置き換えると、ほぼ同じ時間がかかり、「THIS LINE」がないと33ms、あると297msです。データをグローバルメモリに保存するのは正常ですか。共有メモリと同じ時間がかかります。それも「コンパイラの最適化」の一部ですか？

スタックオーバーフローに関する他の同様の問題も確認しました。つまり、データを計算しても保存しないのは無意味であるため、コンパイラの最適化が原因である可能性がある共有メモリにデータを保存するかどうかの間に大きな時間差があるため、コンパイラはそれらの無意味なコードを「削除」しただけです。

行がゲームを変更するのは仮説であるため、同じ理由を共有しているかどうかはわかりません- 「この行」、コメントすると、変数「カウント」がすべての反復で増加し、コメントを外すと、 t は意味があります。

何か案は？お願いします...

4

1 に答える 1

1

多くの場合、比較的小さなコード変更 (カーネル内のコード行の追加または削除など) の結果として大きなパフォーマンスの変化が見られる場合、パフォーマンスの変化はそのコード行の実際のパフォーマンスへの影響によるものではなく、コンパイラがさまざまな最適化の決定を下すため、カーネル内のマシンコードの大規模な追加または削除が発生する可能性があります。

これを確認する比較的簡単な方法は、生成されたマシン コードを確認することです。たとえば、生成されたマシン コードのサイズがソース コードの 1 行の追加または削除によって大幅に変化する場合、コンパイラがコードに大きな影響を与える最適化の決定を行った可能性があります。

これはマシン コードではありませんが、これらの目的では、コンパイラが作成する中間コードである、生成された PTX コードを調べることが合理的な代用となります。

-ptxコンパイル コマンドにスイッチを追加するだけで、ptx を生成できます。

nvcc -ptx mycode.cu

mycode.ptxこれにより、検査可能なファイルが生成されます。当然、通常のコンパイル コマンドに追加のスイッチが必要な場合 (例: -I/path/to/include/files)、このコマンドにも同じスイッチが必要になる場合があります。nvcc マニュアルにはコード生成オプションの詳細が記載されており、PTX について学習するのに役立つ PTXマニュアル.ptxがありますが、生成された PTX のサイズ (たとえば、ファイルの行数) に基づいて大まかなアイデアを得ることができる場合があります。 ）。

于 2013-11-16T13:57:02.457 に答える