2

Fermi ボードを使用して NVidia ビジュアル プロファイラー (CUDA 5.0 ベータ リリースの Eclipse ベースのバージョン) を使用していますが、次の 2 つのパフォーマンス メトリックについて不明な点があります。

  • グローバル ロード/ストア効率は、要求されたトランザクション数に対する実際のメモリ トランザクション数の比率を表します。

  • グローバル メモリ命令リプレイ。最適ではないメモリ結合によって発生したリプレイが原因で発行された命令の割合を表します。

ロード/ストアの効率が 100% (つまり、完全な合体) の場合、グローバル メモリ命令のリプレイは 0 になるべきだという印象を受けました。どうして ？

どうも

4

2 に答える 2

2

私の知る限り、グローバルなロード/ストアの効率はグローバルなメモリアクセスパターンによって決まりますが、グローバルなメモリ命令の再生は主にブランチの分岐によって引き起こされます。したがって、すべてのメモリアクセスが合体していても、ある程度の相違が存在する場合でも、説明したケースが発生する可能性があります。

PS次善のメモリ合体アクセスがグローバルメモリ命令の再生を引き起こす例をいくつか挙げてください。

于 2012-12-12T13:16:03.417 に答える
2

簡単に言えば、単一のワープ トランザクションのサイズには 128 B の制限があるということです (バス幅が原因だと思います)。したがって、ワープに 256 B の合体データが必要な場合は、2 番目の 128 B の命令を再生する必要があります。

一般に、トランザクションは 32B、64B、および 128B セグメントのデータのみを移動します。ワープ トランザクションがこれらのいずれにも適合しない場合は、命令を少なくとも 1 回再生します。合体パターンはこれを避けることはできませんが、トランザクションを最小限に抑えるのに役立ちます。たとえば、ワープ内の Bytes の合体アクセスは、32B トランザクションを取得します。ワープ内で結合された 4B アクセス (int または float) は、単一の 128B トランザクションを取得します。

次のカーネルを検討してください。

__global__ void
gmemtest(const double* const src, double* const dest, const int size,
         const int eleMoved){

  int block_fst = blockIdx.x*blockDim.x*eleMoved;
  size_t thread_fst = block_fst + threadIdx.x*eleMoved;


  #pragma unroll
  for(size_t i = 0; i < eleMoved; i++){
    if( thread_fst + i < size )
      dest[thread_fst + i] = src[thread_fst + i];
  }

elemoved次に、サイズ 1、2、4、および 8 で実行します。カーネルelemovedが大きくなるにつれて、リプレイが増加することがわかります。次のホスト側ループは、128 と 256 のブロック サイズでそれらすべてにヒットします。

  for(size_t j = 1; j<3; j++){

    for(size_t  i = 1; i<=8; i *= 2){

      size_t n_threads = j*128;
      size_t ele_per_thread = i;

      size_t tot_threads = ((SIZE-1)/ele_per_thread)+1;
      size_t n_blocks = ((tot_threads - 1)/n_threads)+1;

      gmemtest<<<n_blocks,n_threads>>>(d_src,d_dest,SIZE,ele_per_thread);
    }
  }

実行nvprof --print-gpu-trace --metrics inst_replay_overheadすると次のことがわかります。

    ==22053== Profiling result:
    Device         Context  Stream   Kernel           Instruction Replay Overhead

   Tesla K20c (0)     1       2    gmemtest(double cons      0.191697
   Tesla K20c (0)     1       2    gmemtest(double cons      0.866548
   Tesla K20c (0)     1       2    gmemtest(double cons      3.472359
   Tesla K20c (0)     1       2    gmemtest(double cons      7.444514
   Tesla K20c (0)     1       2    gmemtest(double cons      0.175090
   Tesla K20c (0)     1       2    gmemtest(double cons      0.912531
   Tesla K20c (0)     1       2    gmemtest(double cons      4.067719
   Tesla K20c (0)     1       2    gmemtest(double cons      7.576686

実際には、ワープに値するデータのようなものを移動している場合、これに遭遇する可能性がありdouble2ます。

本当にパフォーマンス関連の問題に取り組みたい場合は、この講演を十分にお勧めできません: Micikevicius - "Performance Optimization: Programming Guidelines and GPU Architecture Details Behind They"

于 2013-10-07T20:37:18.517 に答える