0

カーネル実行全体のスレッド実行数を徐々にカウントしたい。これに対するネイティブカウンターはありますか、それともそれを行う他の方法はありますか? グローバルメモリ内の変数はスレッドによる同期アクセスを保証しないため、グローバル変数を保持して各スレッドでインクリメントすることは機能しないことを知っています。

4

3 に答える 3

8

スレッド レベルの実行効率を測定するには、さまざまな方法があります。この回答は、さまざまな収集メカニズムのリストを提供します。Robert Crovella の回答は、正確な情報収集を可能にする手動計測方法を提供します。同様の手法を使用して、カーネル内の分岐情報を収集できます。

実行のために起動されたスレッドの数 (静的)

gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

起動されたスレッドの数

gridDim.x * gridDim.y * gridDim.z * ROUNDUP((blockDim.x * blockDim.y * blockDim.z), WARP_SIZE)

この数には、ワープの存続期間中非アクティブなスレッドが含まれます。

これは、PM カウンターの threads_launched を使用して収集できます。

実行されたワープ命令

inst_executed カウンターは、実行/リタイアしたワープ命令の数をカウントします。

ワープ指示が出ました

カウンタ inst_issued は、発行された命令の数をカウントします。inst_issued >= inst_executed. 一部の命令は、狭い実行ユニットへのディスパッチを処理するため、または共有メモリおよび L1 操作でのアドレス分岐を処理するために、実行される命令ごとに複数回発行されます。

実行されたスレッド命令

カウンタ thread_inst_executed は、実行されたスレッド命令の数をカウントします。メトリクス avg_threads_executed_per_instruction は、thread_inst_executed / inst_executed を使用して導出できます。このカウンターの最大値は WARP_SIZE です。

予測されていないオフスレッド命令が実行されました

コンピューティング機能 2.0 以降のデバイスでは、分岐命令の短いシーケンスのパフォーマンスを最適化するために、命令述語を使用してワープ内のスレッドのライトバックを無効にします。

カウンター not_predicated_off_thread_inst_executed は、すべてのスレッドによって実行された命令の数をカウントします。このカウンターは、コンピューティング機能 3.0 以降のデバイスでのみ使用できます。

not_predicated_off_thread_inst_executed <= thread_inst_executed <= WARP_SIZE * inst_executed

この関係は、thread_inst_executed および not_predicated_off_thread_inst_executed カウンターの小さなバグにより、一部のチップではわずかにずれます。

プロファイラ

Nsight Visual Studio Edition 2.x は、前述のカウンターの収集をサポートしています。

Nsight VSE 3.0 は、SASS 命令ごとの統計を収集し、データを表形式で、または高レベルのソース、PTX、または SASS コードの横に表示できる新しい命令カウント実験をサポートしています。情報は、SASS から高レベルのソースにまとめられます。ロールアップの品質は、高品質のシンボル情報を出力するコンパイラの能力に依存します。常にソースと SASS の両方を同時に見ることをお勧めします。この実験では、次の命令ごとの統計を収集できます。

a. inst_executed b. thread_inst_executed (またはアクティブなマスク) c. ｎｏｔ＿ｐｒｅｄｉｃａｔｅｄ＿ｏｆｆ＿ｔｈｒｅａｄ＿ｉｎｓｔ＿ｅｘｅｃｕｔｅｄ（アクティブな述語マスク）ｄ．active_mask のヒストグラム e. predicate_mask のヒストグラム

Visual Profiler 5.0 は、前述の SM カウンターを正確に収集できます。nvprof は、SM ごとの詳細を収集して表示できます。Visual Profiler 5.x は、Nsight VSE 3.0 で使用可能な命令ごとの統計の収集をサポートしていません。古いバージョンの Visual Profiler および CUDA コマンド ライン プロファイラーは、前述のカウンターの多くを収集できますが、結果はツールの 5.0 以降のバージョンほど正確ではない場合があります。

于 2013-05-29T02:08:25.433 に答える
3

多分このようなもの：

__global__ void mykernel(int *current_thread_count, ...){
  atomicAdd(current_thread_count, 1);
  // the rest of your kernel code
  }


int main() {
  int tally, *dev_tally;
  cudaMalloc((void **)&dev_tally, sizeof(int));
  tally = 0;
  cudaMemcpy(dev_tally, &tally, sizeof(int), cudaMemcpyHostToDevice);
  ....
  // set up block and grid dimensions, etc.
  dim3 grid(...);
  dim3 block(...)

  mykernel<<<grid, block>>>(dev_tally, ...);
  cudaMemcpy(&tally, dev_tally, sizeof(int), cudaMemcpyDeviceToHost); 
  printf("total number of threads that executed was: %d\n", tally);
  ....

  return 0;
  }

アトミック関数の詳細については、こちらをご覧ください

コメントで多くの人が表明した混乱の理由の 1 つは、mykernel(正常に実行されたと仮定して) が完了すると、誰もtallyが最終的に次の値になることを期待しているからです。grid.x*grid.y*grid.z*block.x*block.y*block.z

于 2013-05-28T15:48:20.097 に答える
0

特定のパス ブランチのスレッド数を計算する方法はないと思います。たとえば、ヒストグラムの場合、次のようにすると便利です。

PS: ヒストグラムは、各色のピクセルを数えることに関するものです。

for (i=0; i<256; i++) // 256 colors, 1 pixel = 1 thread if (threadidx.x == i) Histogramme[i] = CUDA_NbActiveThreadsInBranch() // Threads having i as color

于 2015-03-17T18:40:16.753 に答える