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

Question

我正在使用 NVidia 视觉分析器(来自 CUDA 5.0 beta 版本的基于 eclipse 的版本)和 Fermi 板，我不了解其中两个性能指标:

全局加载/存储效率表示实际内存事务数与请求事务数的比率。

全局内存指令重放，表示由于次优内存合并导致的重放而发出的指令的百分比。

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

谢谢

Answer 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];
  }

现在用 elemoved 运行它大小为 1、2、4 和 8。您会发现内核的重放增加为 elemoved。变大。以下主机端循环将以 128 和 256 的 block 大小命中所有这些。

  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

在实践中，如果您要移动值(value)为 double2 的经线之类的东西，您可能会遇到这种情况。数据。

如果你真的想讨论与性能相关的问题，我不能推荐这个话题: Micikevicius - "Performance Optimization: Programming Guidelines and GPU Architecture Details Behind Them"