CUDA同步和读取全局内存

Question

我有一个类似于这样的东西:

__global__ void globFunction(int *arr, int N) {
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;
    // calculating and Writing results to arr ...
    __syncthreads();
    // reading values of another threads(ex i+1)
    int val = arr[idx+1]; // IT IS GIVING OLD VALUE
}


int main() {
    // declare array, alloc memory, copy memory, etc.
    globFunction<<< 4000, 256>>>(arr, N); 
    // do something ...
    return 0;
}

为什么我读取arr[idx+1]的时候得到的是旧值？我已经调用了__syncthreads，所以我期望看到更新后的值。我做错了什么？我是在读取缓存还是其他地方？

- nosbor

---

__syncthreads是一个块级别的同步原语，而不是网格级别的。根据你最近的几个问题，看起来你应该花一些时间阅读CUDA文档。 - talonmies

好的，那么它是不可能读取的吗？ - nosbor

@sidyll：这是CUDA的C语言代码，请保留C标签。 - Kerrek SB

Answer 1

使用__syncthreads()函数只同步当前块中的线程。在这种情况下，当您启动内核时，每个块中有256个线程。因此，在给定的数组中，对于每个跨越到另一个线程块的索引值，您将会读取一个与当前块中的线程不同步的全局内存值。

为了解决这个问题，您可以使用__shared__ CUDA指令创建共享的线程本地存储器，允许块中的线程彼此共享信息，但防止来自其他块的线程访问为当前块分配的内存。一旦块内的计算完成（您可以使用__syncthreads()执行此任务），您就可以将共享块级存储器中的值复制回全局可访问的内存中。

您的内核可能如下所示：

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

如果您必须在块之间同步线程，则应寻找另一种解决问题的方式，因为CUDA编程模型在问题可以分解为块并且仅需要在块内进行线程同步时最有效。

- Jason

---

警告，这是危险代码。编译器可能会重新排列对标志和数组的写入，导致竞争条件。为了正确，您可能需要在此处添加__threadfence()。一般来说，没有原子操作的块间通信必须小心处理，如果可能的话最好找另一种方法... - harrism

感谢提供关于__threadfence的信息。我已将其添加到代码中，并根据Nvidia文档，应该可以正常工作。由于每个线程只写入一个标志（每个线程都有一个标志），因此我认为不需要原子操作。 - Jason

@Jason: 你的第二段代码假设线程idx和线程idx+1同时被调度和运行。然而，CUDA执行模型没有做出这方面的保证。抱歉，但是这段代码是有问题的，它很容易在演示条件下死锁。 - talonmies

好的，我没有意识到关于CUDA运行时的问题。我已经修改了我的答案以反映这一点。如果您发现有什么不对的地方，我会非常感激您的进一步评论。 - Jason