cuda - __syncthreads() 是否同步网格中的所有线程？

Question

__syncthreads() 是同步网格中的所有线程还是仅同步当前扭曲或 block 中的线程？

此外，当特定 block 中的线程遇到(在内核中)以下行时

__shared__  float srdMem[128];

他们只会声明这个空间一次(每个 block )吗？

它们显然都是异步操作的，因此如果 block 22 中的线程 23 是第一个到达该行的线程，然后 block 22 中的线程 69 是最后一个到达该行的线程，则线程 69 将知道它已经被声明？

Answer 1

__syncthreads()命令是一个 block 级同步屏障。这意味着当 block 中的所有线程都到达屏障时，可以安全地使用它。也可以使用__syncthreads()在条件代码中，但仅当所有线程对此类代码进行相同的评估时，否则执行可能会挂起或产生意外的副作用 [4] .

使用__syncthreads()的示例:(source)

__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;        
}

要同步网格中的所有线程，当前没有 native API 调用。在网格级别上同步线程的一种方法是使用连续的内核调用，因为此时所有线程都会从同一点结束并重新开始。它通常也称为 CPU 同步或隐式同步。因此它们都是同步的。

使用此技术的示例 ( source ):

关于第二问题。 是的，它确实声明了每个 block 指定的共享内存量。请考虑到可用共享内存的数量是按SM测量的。因此，人们应该非常小心如何使用共享内存以及启动配置。