CUDA同步和读取全局内存

2023年3月25日 319次阅读

我有这样的事情：

__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,所以我希望看到更新的值.我做错了什么？我在读缓存还是什么？

最佳答案 使用__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编程模型最有效,并且线程同步只需要在块内进行.

点赞