0

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

__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 并且时间保持不变。如果我大大减小循环大小也是一样的。为什么会这样？

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

4

2 回答 2

4

如前所述，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，就不会执行存储，因此不会影响循环的时间。

于 2013-08-30T05:02:12.167 回答
2

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

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

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

于 2013-08-29T22:54:35.187 回答