2

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

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

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

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

谢谢

4

2 回答 2

2

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

PS您能否举一些示例，其中次优内存合并访问会导致全局内存指令重播？

于 2012-12-12T13:16:03.417 回答
2

简短的回答是，单个 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 的块大小命中所有这些。

  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 - “性能优化：背后的编程指南和 GPU 架构细节”

于 2013-10-07T20:37:18.517 回答