cuda - GPGPU-CUDA : global store efficiency

Question

我试图根据 NVidia 分析器的“global store efficiency”值，弄清楚我的一个内核的全局内存写访问的合并情况(我在 Fermi GPU 上使用 CUDA 5 工具包预览版)。

据我所知，这个值是请求的内存事务与实际执行的事务 nb 的比率，因此反射(reflect)了访问是否完全合并(100% 效率)。

现在，对于 32 的线程 block 宽度，并将浮点值作为输入和输出，以下测试内核如预期的那样为全局加载和全局存储提供 100% 的效率:

__global__ void dummyKernel(float*output,float* input,size_t pitch)
{
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  int offset = y*pitch+x;
  float tmp = input[offset];
  output[offset] = tmp;
}

我不明白的是，为什么当我开始在输入读取和输出写入之间添加有用的代码时，全局存储效率开始下降，而我没有更改内存写入模式或线程 block 几何结构？不过，如我所料，全局负载保持在 100%。

有人能解释一下为什么会这样吗？我想，由于给定 warp 中的所有 32 个线程同时(根据定义)执行输出存储指令并使用“合并友好”模式，无论我以前做什么，我仍然应该得到 100%，但显然我一定是误解了一些东西全局存储效率的含义，或者全局存储合并的条件。

谢谢，

编辑:

这是一个例子:如果我使用这段代码(只是在输入上添加一个“round”操作)，全局存储效率会从 100% 下降到 95%

__global__ void dummyKernel(float*output,float* input,size_t pitch)
{
  unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
  unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
  int offset = y*pitch+x;
  float tmp = round(input[offset]);
  output[offset] = tmp;
}

Answer 1

不确定是否是这种情况，但 round 可能会将其参数转换为 double ，如果存在寄存器溢出，则每个线程将访问 8 字节的内存，然后将其强制转换为 4 字节的 tmp。访问 8 个字节会将合并减少到半扭曲。

但是，我相信寄存器溢出不应该发生，因为内核中的局部变量数量很少。您可以使用 nvcc --ptxas-options=-v 检查溢出情况。