cuda控制分歧

Question

说我有3个共享内存阵列:ab c.我不确定跟随线程安排是否会导致控制分歧,

if (threadIdx < 64)
{
    if (threadIdx == 1)
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*a[threadIdx];
    else
        for (int i = 0; i < N; i++)
            c += a[threadIdx]*b[threadIdx];
}

如果确实如此,它对性能的影响有多大？有没有有效的方法来处理这个问题？谢谢

Answer 1

根据块的尺寸,第一个条件threadIdx.x < 64(注意.x)可能不会导致任何分歧.例如,如果您有一个具有维度的块,(128,1,1)那么前两个warp(以锁定步骤执行的32个线程组)将进入该if块,而最后两个将绕过它.由于整个经线是这样或那样的,所以没有分歧.

有条件的threadIdx.x == 1会导致分歧,但它会有非常适度的成本.实际上,在许多情况下,CUDA将能够使用单个指令实现条件表达式.例如,像min,max和这样的操作abs通常用单个指令实现,并且绝对不会产生分歧.您可以在PTX手册中阅读有关此类说明的信息.

一般来说,你不应该过分关注如上所述的适度控制流量分歧.在诸如此类的情况下,分歧会让你陷入困境

if (threadIdx.x % 4 == 0)
  // do expensive operation
else if (threadIdx.x % 4 == 1)
  // do expensive operation
else if (threadIdx.x % 4 == 2)
  // do expensive operation
else
  // do expensive operation

其中"昂贵的操作"是需要10或100条指令的操作.在这种情况下,由if陈述引起的分歧会使效率降低75%.

请记住,线程分歧比(1)高级算法选择和(2)内存局部性/合并更少关注.很少有CUDA程序员应该关注你的例子中的那种分歧.

Answer 2

如果每个块有多个线程,我会期望每个块的一个warp中的分歧(无论哪个块保存线程1).

但是,两个循环之间的区别仅在于要访问的内存,而不是在指令中.所以,我会这样做......

if (threadIdx.x < 64)
{
    //this conditional might diverge
    if (threadIdx.x == 1)
        ptr = a;
    else
        ptr = b;

    //but obviously this part will not
    for (int i = 0; i < N; i++)
        c += a[threadIdx]*ptr[threadIdx];
}