3

我了解如何#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 万个循环会导致寄存器压力，我不确定编译器是否会这样做），如果有的话，什么是“公平”数字？我如何确定这个限制？

4

2 回答 2

3

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

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

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没有展开。如果您比我更有耐心，您可能可以使用模板来找到此代码的确切编译器限制，尽管我怀疑这对了解特别有指导意义。

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

于 2015-12-30T10:19:41.793 回答
1

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

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

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

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

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

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

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

于 2015-12-30T23:35:28.003 回答