我了解其#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万个循环将导致寄存器压力，并且我不确定编译器是否会进行展开），那么“公平”的数字是多少（如果有）？以及如何确定该限制？
您的示例内核是完全串行的，无论如何在现实世界中都不是循环展开的有用用例，但让我们将自己限制在编译器将执行多少次循环展开的问题上。
这是内核的可编译版本，带有一些模板修饰：
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部署多达10个内核。该LIMIT=8192案件尚未解开。如果您有更多的耐心，可以尝试使用此模板为该代码找到确切的编译器限制，尽管我怀疑这样做特别有启发性。
您还可以通过编译器自己查看所有严重展开的版本都使用相同数量的寄存器（由于内核的琐碎性质）。
本文收集自互联网，转载请注明来源。
如有侵权，请联系[email protected] 删除。
我来说两句