3

NVidiaのプロファイラーの「グローバルストア効率」の値に基づいて、カーネルの1つのグローバルメモリ書き込みアクセスがどれだけうまく合体しているかを把握しようとしています（Fermi GPUでCUDA5ツールキットプレビューリリースを使用しています）。

私が理解している限り、この値は、実行されたトランザクションの実際のnbに対する要求されたメモリトランザクションの比率であり、したがって、アクセスがすべて完全に合体されている（100％の効率）かどうかを反映しています。

ここで、スレッドブロック幅が32で、float値を入力と出力として使用する場合、次のテストカーネルは、予想どおり、グローバルロードとグローバルストアの両方で100％の効率を提供します。

__global__ void dummyKernel(float*output,float* input,size_t pitch)
{
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  int offset = y*pitch+x;
  float tmp = input[offset];
  output[offset] = tmp;
}

私が理解していないのは、入力読み取りと出力書き込みの間に有用なコードを追加し始めると、メモリ書き込みパターンやスレッドブロックジオメトリを変更していないのに、グローバルストアの効率が低下し始める理由です。ただし、予想どおり、グローバル負荷は100％のままです。

誰かがなぜこれが起こるのかを明らかにしてくれませんか？特定のワープの32スレッドすべてが（定義上）出力ストア命令を同時に実行し、「合体しやすい」パターンを使用しているので、以前と同じように100％取得する必要があると思いましたが、明らかに何かを誤解している必要があります。グローバルストアの効率の意味、またはグローバルストアの合体の条件のいずれか。

どうも、

編集 ：

次に例を示します。このコードを使用すると（入力に「ラウンド」操作を追加するだけ）、グローバルストアの効率が100％から95％に低下します。

__global__ void dummyKernel(float*output,float* input,size_t pitch)
{
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  int offset = y*pitch+x;
  float tmp = round(input[offset]);
  output[offset] = tmp;
}
4

2 に答える 2

0

これが当てはまるかどうかはわかりませんが、roundはおそらく引数をdoubleに変換し、レジスタがスピルすると、各スレッドは8バイトのメモリにアクセスし、4バイトのtmpに強制変換されます。8バイトにアクセスすると、合体が半分に減少します。

ただし、カーネル内のローカル変数の数が少ないため、レジスタのスピルは発生しないはずです。流出については、nvcc --ptxas-options=-vで確認できます。

于 2012-06-28T15:11:52.027 に答える
0

残念ながら、問題が見つかりました。この単純なテストコードをデバッグモードでプロファイリングしていました。これにより、ほとんどのメトリックで完全にワイルドな数値が得られます。リリースモードでの再プロファイリングにより、期待される結果が得られました。どちらの場合も、ストアの効率は100％です。

于 2012-06-29T09:55:35.027 に答える