CUDA shuffle的竞争条件？

Question

使用shuffle命令，当两个不同的线程并发地尝试更新相同的寄存器值时，是否存在竞争条件/丢失的更新？

Answer 1

这是此处提供的较晚的答案，用于从未回答列表中删除此问题。

来自CUDA C编程指南

__shfl()内部函数允许在不使用共享内存的情况下在warp内的线程之间交换变量

其思想是线程i可以读取分配给线程j的寄存器r的值，但不能更改该值。因此，正如上面的评论所指出的，没有竞争条件。

CUDA C编程指南》还提供了以下跨偏差广播单个值的示例

global__ void bcast(int arg) {

    int laneId = threadIdx.x & 0x1f;
    int value;

    if (laneId == 0)          // Note unused variable for
    value = arg;              // all threads except lane 0
    value = __shfl(value, 0); // Get "value" from lane 0
    if (value != arg) printf("Thread %d failed.\n", threadIdx.x); }

void main() {
    bcast<<< 1, 32 >>>(1234);
    cudaDeviceSynchronize();
}

在此示例中，分配给warp中的线程0的value寄存器的值被广播到warp中的所有其他线程，并分配给本地value寄存器。所有其他线程不会尝试(，但也不能)更改分配给线程0的value寄存器的值。