确定 CUDA 中#pragma unroll N 的最佳值

Question

我了解#pragma unroll 的工作原理，但如果我有以下示例：

__global__ void
test_kernel( const float* B, const float* C, float* A_out)
{
  int j = threadIdx.x + blockIdx.x * blockDim.x;
  if (j < array_size) {
     #pragma unroll
     for (int i = 0; i < LIMIT; i++) {
       A_out[i] = B[i] + C[i];
     }
  }
}

我想确定上面内核中LIMIT 的最佳值，它将以x 线程数和y 块数启动。 LIMIT 可以是从 2 到 1<<20 的任意位置。由于 100 万对于变量来说似乎是一个非常大的数字（展开的 100 万个循环会导致寄存器压力，我不确定编译器是否会这样做），如果有的话，什么是“公平”数字？我如何确定这个限制？

Answer 1

您的示例内核是完全串行的，无论如何都不是循环展开的有用的现实世界用例，但让我们将自己限制在编译器将执行多少循环展开的问题。

这是您的内核的可编译版本，带有一些模板装饰：

template<int LIMIT>
__global__ void
test_kernel( const float* B, const float* C, float* A_out, int array_size)
{
  int j = threadIdx.x + blockIdx.x * blockDim.x;
  if (j < array_size) {
     #pragma unroll
     for (int i = 0; i < LIMIT; i++) {
       A_out[i] = B[i] + C[i];
     }
  }
}

template __global__ void test_kernel<4>(const float*, const float*, float*, int);
template __global__ void test_kernel<64>(const float*, const float*, float*, int);
template __global__ void test_kernel<256>(const float*, const float*, float*, int);
template __global__ void test_kernel<1024>(const float*, const float*, float*, int);
template __global__ void test_kernel<4096>(const float*, const float*, float*, int);
template __global__ void test_kernel<8192>(const float*, const float*, float*, int);

您可以将其编译为 PTX 并亲自查看（至少使用 CUDA 7 版本编译器和默认计算能力 2.0 目标架构），具有高达 LIMIT=4096 的内核已完全展开。 LIMIT=8192 案例未展开。如果您比我更有耐心，您可能可以使用模板来找到此代码的确切编译器限制，尽管我怀疑这对了解特别有启发性。

您还可以通过编译器亲眼看到所有大量展开的版本都使用相同数量的寄存器（因为您的内核很简单）。

Answer 2

CUDA 利用线程级并行性（通过将工作拆分为多个线程来公开）和指令级并行性（CUDA 通过在编译代码中搜索独立指令来发现）。

@talonmies 的结果显示，您的循环可能在 4096 到 8192 次迭代之间展开，这让我感到惊讶，因为循环展开在现代 CPU 上的回报急剧减少，其中大多数迭代开销已通过分支等技术优化掉预测和推测执行。

在 CPU 上，我怀疑展开超过 10 到 20 次迭代会带来很多好处，并且展开的循环会在指令缓存中占用更多空间，因此展开也会产生成本。 CUDA 编译器在确定展开多少时将考虑成本/收益权衡。所以问题是，展开 4096+ 次迭代可能有什么好处？我认为这可能是因为它为 GPU 提供了更多代码，它可以在其中搜索独立指令，然后使用指令级并行性并发运行。

循环的主体是A_out[i] = B[i] + C[i];。由于循环中的逻辑不访问外部变量，也不访问循环早期迭代的结果，因此每次迭代都独立于所有其他迭代。所以i 不必依次增加。即使循环以完全随机的顺序在0 和LIMIT - 1 之间迭代i 的每个值，最终结果也会相同。该属性使循环成为并行优化的良好候选者。

但有一个问题，这就是我在评论中提到的。只有当A 缓冲区与B 和C 缓冲区分开存储时，循环的迭代才是独立的。如果您的A 缓冲区与内存中的B 和/或C 缓冲区部分或完全重叠，则会创建不同迭代之间的连接。一次迭代现在可以通过写入A 来更改另一次迭代的B 和C 输入值。因此，根据两个迭代中的哪一个先运行，您会得到不同的结果。

多个指向内存中相同位置的指针称为指针别名。因此，一般来说，指针别名会导致看似独立的代码段之间的“隐藏”连接，因为一段代码通过一个指针完成的写入可能会改变另一段代码从另一个指针读取的值。默认情况下，CPU 编译器生成的代码考虑了可能的指针别名，生成的代码无论如何都会产生正确的结果。问题是 CUDA 做了什么，因为回到 talonmies 的测试结果，我能看到如此大量展开的唯一原因是它为指令级并行性打开了代码。但这意味着 CUDA 在这种特殊情况下不考虑指针别名。

回复。关于运行多个线程的问题，当您增加线程数时，常规串行程序不会自动变为并行程序。您必须确定可以并行运行的工作部分，然后在您的 CUDA 内核中表达出来。这就是所谓的线程级并行性，它是提高代码性能的主要来源。此外，CUDA 会在每个内核中搜索独立的指令，并可能同时运行这些指令，这就是指令级并行。高级 CUDA 程序员可能会牢记指令级并行性并编写有助于实现这一点的代码，但我们凡人应该只关注线程级并行性。这意味着您应该再次查看您的代码并考虑可能能够并行运行。由于我们已经得出结论，您的循环体是并行化的良好候选者，因此您的工作就是重写内核中的串行循环，以向 CUDA 表达如何并行运行单独的迭代。