循环的cuda测量答案

【问题标题】：Cuda measurement of loop循环的cuda测量
【发布时间】：2013-08-30 13:02:12
【问题描述】：

我在 CUDA Fermi GPU 上启动了一个非常简单的内核 >>。

__global__ void kernel(){
int x1,x2;

x1=5;
x2=1;

for (int k=0;k<=1000000;k++)
  {
   x1+=x2;

  }
}

内核非常简单，它执行 10^6 次加法，并且不会将任何内容传回全局内存。结果是正确的，即循环 x1（在其所有 512 个线程实例中）包含 10^6 + 5

我正在尝试测量内核的执行时间。同时使用 Visual Studio 并行 nsight 和 nvvp。 Nsight 为 2.5 微秒，nvvp 为 4 微秒。

问题如下：我可能会大幅增加循环的大小，例如增加到 10^8 并且时间保持不变。如果我大大减小循环大小也是一样的。为什么会这样？

请注意，如果我在循环内使用共享内存或全局内存，则测量值会反映正在执行的工作（即存在比例性）。

【问题讨论】：

标签： cuda measurement

【解决方案1】：

如前所述，CUDA 编译器优化在删除死代码方面非常积极。因为x2 不参与写入内存的值，所以可以删除它和循环。编译器还会预先计算任何可以在编译时推导出的结果，因此如果编译器知道循环中的所有常量，它可以计算最终结果并将其替换为常量。

要解决这两个问题，请像这样重写您的代码：

__global__ 
void kernel(int *out, int x0, bool flag)
{
    int x1 = x0, x2 = 1;

    for (int k=0; k<=1000000; k++) {
       x1+=x2;
    }

    if (flag) out[threadIdx.x + blockIdx.x*blockDim.x] = x1;
}

然后像这样运行它：

kernel<<<1,512>>>((int *)0, 5, false);

通过将x1 的初始值作为参数传递给内核，您可以确保编译器无法获得循环结果。该标志使内存存储有条件，然后内存存储使整个计算无法安全删除。只要在运行时将标志设置为 false，就不会执行存储，因此不会影响循环的时间。

【讨论】：

【解决方案2】：

因为编译器消除了死路。你的代码实际上并没有做任何事情。看组装。

如果您确实看到了该值，那么编译器可能刚刚优化了循环，因为它可以在编译时知道该值。

当您将寄存器内容写入共享内存时，编译器无法保证不会使用结果，因此会实际计算该值。换句话说，您计算的值最终必须在某个地方使用或写入内存，否则它的计算将被丢弃。

【讨论】：