复制到cuda的共享内存

Question

在CUDA编程中,如果我们想要使用共享内存,我们需要将数据从全局内存带到共享内存.线程用于传输此类数据.

我在某处(在线资源中)读到最好不要将块中的所有线程都用于将数据从全局内存复制到共享内存.这样的想法是有道理的,所有线程不会一起执行.warp中的线程一起执行.但我担心的是所有的经线都没有按顺序执行.比如说,一个带有线程的块被分为3个warp:war p0(0-31个线程),warp 1(32-63个线程),warp 2(64-95个线程).不保证warp 0将首先执行(我是对的吗？).

那么我应该使用哪些线程将数据从全局复制到共享内存？

Answer 1

要使用单个warp加载共享内存数组,只需执行以下操作:

__global__
void kernel(float *in_data)
{
    __shared__ float buffer[1024];

    if (threadIdx.x < warpSize) {
        for(int i = threadIdx; i  <1024; i += warpSize) {
            buffer[i] = in_data[i];
        }
    }
    __syncthreads();

    // rest of kernel follows
}

[免责声明:用浏览器编写,从未测试过,使用风险自负]

这里的关键点是使用__syncthreads()确保块中的所有线程等到执行加载到共享内存的warp已完成加载.我发布的代码使用了第一个warp,但您可以通过将块内的线程索引除以warpSize来计算warp数.我还假设了一维块,在2D或3D块中计算线程索引是微不足道的,所以我把它作为练习留给读者.

Answer 2

当块被分配给多处理器时，它会驻留在那里，直到该块内的所有线程都完成为止，在此期间，warp 调度程序在具有就绪操作数的 warp 之间进行混合。因此，如果多处理器上有一个块具有三个 warp，并且只有一个 warp 正在从全局内存获取数据到共享内存，而其他两个 warp 保持空闲状态并且可能正在等待__syncthreads()屏障，那么您不会丢失任何内容，并且仅受到全局内存延迟的限制无论如何你都会是的。一旦抓取完成，扭曲就可以继续进行它们的工作。

因此，不需要保证 warp0 首先执行，您可以使用任何线程。唯一需要记住的两件事是确保尽可能多地联合访问全局内存并避免银行冲突。