我了解 #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之间任何位置。1百万似乎对于变量来说是非常大的数字（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);
LIMIT = 4096的内核完全展开。 LIMIT = 8192 情况未展开。如果你比我更有耐心，可以尝试使用模板来查找此代码的确切编译器限制，但我认为了解这一点并不特别有教育意义。A_out[i] = B[i] + C[i];。由于循环中的逻辑没有访问外部变量，也没有访问之前迭代的结果，因此每次迭代都是独立的。所以i不必按顺序递增。即使循环以完全随机的顺序迭代i的每个值，最终结果也是相同的。这种特性使得循环成为并行优化的好选择。A缓冲区与B和/或C缓冲区分开存储时，循环的迭代才是独立的。如果A缓冲区部分或全部重叠B和/或C缓冲区，则会创建不同迭代之间的连接。一次迭代可能通过写入A来更改另一次迭代的B和C输入值。因此，取决于哪个迭代先运行，您将获得不同的结果。
A、B和C是什么，它们存储在哪里？为什么循环完全串行？您希望从展开完全串行的循环（看起来像线程本地变量）中获得什么优势？ - talonmiestest_kernel<<<1, 1>>>(d_idata_B, d_idataC, d_odataA);的方式进行调用。 - Blizzard