cuda - #pragma unroll 到底有什么作用？对线程数有影响吗？

Question

我是 CUDA 新手，我无法理解循环展开。我编写了一段代码来理解该技术

__global__ void kernel(float *b, int size)
{
    int tid = blockDim.x * blockIdx.x + threadIdx.x;
 #pragma unroll
    for(int i=0;i<size;i++)
        b[i]=i;
}

上面是我的核函数。在 main 中，我将其称为如下

int main()
{
    float * a; //host array
    float * b; //device array
    int size=100;

    a=(float*)malloc(size*sizeof(float));
    cudaMalloc((float**)&b,size);
    cudaMemcpy(b, a, size, cudaMemcpyHostToDevice);

    kernel<<<1,size>>>(b,size); //size=100

    cudaMemcpy(a, b, size, cudaMemcpyDeviceToHost);

    for(int i=0;i<size;i++)
        cout<<a[i]<<"\t";

    _getch();

    return 0;
}

这是否意味着我有 size*size=10000 个线程正在运行来执行该程序？循环展开时是否创建了 100 个？

Answer 1

没有。这意味着您使用一个 block 调用了 CUDA 内核，并且该 block 有 100 个事件线程。您将 size 作为第二个函数参数传递给内核。在您的内核中，这 100 个线程中的每一个都执行 for 循环 100 次。

#pragma unroll是一种编译器优化，例如可以替换类似

的代码

for ( int i = 0; i < 5; i++ )
    b[i] = i;

与

b[0] = 0;
b[1] = 1;
b[2] = 2;
b[3] = 3;
b[4] = 4;

通过输入 #pragma unroll指令就在循环之前。展开版本的好处是处理器的处理负载较少。如果是for循环版本，处理时，除了分配每个i至b[i] ，涉及i初始化，评估i<5 6次，并递增i 5次。而第二种情况，只涉及备案b数组内容(如果稍后使用 int i=5; 则可能加上 i)。循环展开的另一个好处是增强指令级并行性 (ILP)。在展开版本中，处理器可能会将更多操作插入处理管道，而不必担心 for每次迭代中的循环条件。

类似 this 的帖子解释 CUDA 无法展开运行时循环。在你的情况下，CUDA编译器没有任何线索size将为 100，因此不会发生编译时循环展开，因此如果强制展开，最终可能会损害性能。

如果您确定size对于所有执行都是 100，您可以像下面这样展开循环:

#pragma unroll
for(int i=0;i<SIZE;i++)  //or simply for(int i=0;i<100;i++)
    b[i]=i;

其中SIZE在编译时已知 #define SIZE 100 .

我还建议您在代码中进行适当的 CUDA 错误检查(解释 here )。