3

我正在尝试根据 NVidia 分析器的“全局存储效率”值（我在 Fermi GPU 上使用 CUDA 5 工具包预览版）来确定我的一个内核的全局内存写入访问的合并程度。

据我了解，该值是请求的内存事务与实际执行的事务 nb 的比率，因此反映了访问是否全部完美合并（100% 效率）。

现在，对于 32 的线程块宽度，并将浮点值作为输入和输出，以下测试内核为全局加载和全局存储提供了 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;
}

我不明白为什么当我开始在输入读取和输出写入之间添加有用的代码时，全局存储效率开始下降，而我没有改变内存写入模式或线程块几何？不过，正如我所料，全球负载保持在 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;
}
4

2 回答 2

0

好吧，真丢脸，我发现了问题：我在调试模式下分析这个简单的测试代码，它为大多数指标提供了完全疯狂的数字。在发布模式下重新分析给了我预期的结果：两种情况下的存储效率都是 100%。

于 2012-06-29T09:55:35.027 回答
0

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

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

于 2012-06-28T15:11:52.027 回答