cuda - CUDA shuffle 的竞争条件？

Question

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

Answer 1

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

来自 CUDA C 编程指南

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

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

CUDA C 编程指南还提供了以下示例，以通过 warp 广播单个值

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();
}

在这个例子中，value分配给warp中线程的寄存器的值0被广播到warp中的所有其他线程并分配给本地value寄存器。所有其他线程都没有尝试（但也不能）更改value分配给线程的寄存器的值0。