0

私はCUDAの助けを借りて65536要素の配列を減らしようとしています（その中の要素の合計を計算します）。カーネルは次のようになります（今のところ* dev_distanceFloatsとインデックス引数は無視してください）

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
}

antそれは256スレッドの1つのブロックとして起動しました：

kernel_calcSum <<<1,  256 >>>(dev_spFloats1, dev_distanceFloats, index);

これまでのところ、256個のスレッドのそれぞれがグローバルメモリから256個の要素を取得し、ローカル変数mySumでその合計を計算します。カーネルの実行時間は約45ミリ秒です。次のステップは、ブロック内の256スレッド間に共有メモリを導入することです（mySumの合計を計算するため）。したがって、カーネルは次のようになります。

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  __shared__ float sdata[256];
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
  sdata[tid] = mySum;
}

共有メモリへの書き込みを追加しましたが、実行時間が45ミリ秒から258ミリ秒に増加します（NVidia Visual Profiler 5.0.0を使用してこれをチェックしています）。sdata変数に書き込むときに、スレッドごとに8つのバンク競合があることに気付きました（32バンクの​​機能3.0を持つGTX670を使用しています）。実験として、カーネルを起動するときにスレッドを32に減らしようとしましたが、それでも時間は258ミリ秒です。

質問1：私の場合、共有メモリへの書き込みに時間がかかるのはなぜですか？質問2：「実行計画」（メモリアクセス、競合などのタイミング）を詳細に示すツールはありますか？

あなたの提案をありがとう。

更新： カーネルで遊んでいます-スレッドごとにsdataを定数に設定しました：

__global__ void kernel_calcSum(float *d, float *dev_distanceFloats, int index) {
  __shared__ float sdata[256];
  int tid = threadIdx.x;
  float mySum = 0;
  for (int e = 0; e < 256; e++) {
    mySum += d[tid + e];
  }
  sdata[tid] = 111;
}

タイミングは48ミリ秒に戻ります。したがって、sdata [tid]=mySum;を変更します。to sdata [tid] = 111; これを作りました。

このコンパイラの最適化（この行を削除しただけでしょうか？）または何らかの理由でローカルメモリ（レジスタ？）から共有へのコピーに時間がかかりますか？

4

2 に答える 2

0

共有メモリはこれには適切ではありません。必要なのは、ワープアトミック操作です。ワープ内で合計し、ワープ間の中間結果を伝達します。CUDAでのこの出荷を示すサンプルコードがあります。

要素を合計することは、大規模な並列化があまり役に立たず、実際にはGPUがCPUよりも優れている可能性があるタスクの1つです。

于 2013-03-15T10:37:11.897 に答える
0

どちらのカーネルも、カーネルの終了後もアクセス可能な結果を​​メモリに書き出さないため、何もしません。

最初のケースでは、コンパイラはこれに気づき、計算全体を最適化するのに十分賢いです。共有メモリが関係する2番目のケースでは、共有メモリを介した情報の流れを追跡するのが難しいため、コンパイラはこれに気づきません。したがって、計算はに残されます。

（すでに行っているように）グローバルメモリへのポインタを渡し、このポインタを介して結果を書き出します。

于 2013-03-15T11:06:10.543 に答える