CUDA：加载/存储效率与全局内存指令重放之间的关系答案

【问题标题】：CUDA: relation between load/store efficiency and global memory instruction replayCUDA：加载/存储效率与全局内存指令重放之间的关系
【发布时间】：2012-08-02 12:26:58
【问题描述】：

我正在使用 NVidia 可视化分析器（CUDA 5.0 beta 版本中基于 eclipse 的版本）和 Fermi 板，我对其中两个性能指标有一些不理解：

全局加载/存储效率表示实际内存事务数与请求事务数之比。
全局内存指令重放，表示由于次优内存合并导致重放而发出的指令的百分比。

我的印象是，如果加载/存储效率为 100%（即完美合并），则全局内存指令重播应该为 0，但我已经看到了 100% 效率和非零全局内存指令重播的示例。怎么会？

谢谢

【问题讨论】：

标签： cuda

【解决方案1】：

简短的回答是，单个 warp 事务的大小限制为 128 B（我相信由于总线宽度）。因此，如果您的 warp 需要 256 B 的合并数据，那么您必须重播第二个 128 B 的指令。

一般来说，事务只移动 32B、64B 和 128B 段中的数据。如果您的 warp 交易不适合其中之一，那么您将至少重播一次指令。合并模式无法避免这种情况，但它们确实有助于最小化事务。例如，warp 中 Bytes 的合并访问为您提供 32B 事务。 Warp 中的合并 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];
  }

现在在大小为 1、2、4 和 8 时使用 elemoved 运行它。您会发现内核的重放随着 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 Them"

【讨论】：

【解决方案2】：

据我所知，全局加载/存储效率是由全局内存访问模式决定的，而全局内存指令重放主要是由分支分歧引起的。因此，即使所有内存访问都合并但存在一些分歧，您描述的情况也可能会发生。

附：您能否举一些次优内存合并访问导致全局内存指令重放的例子？

【讨论】：