在 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 万个循环会导致寄存器压力，我不确定编译器是否会展开），什么是 "fair" 数字（如果有）？我如何确定该限制？

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 利用 thread-level 并行性和 instruction-level 并行性，前者通过将工作拆分为多个线程，后者是 CUDA 通过在编译代码中搜索独立指令发现的。

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

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

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

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

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

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

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

Determining the optimal value for #pragma unroll N in CUDA

cuda

pragma

loop-unrolling