复制到cuda中的共享内存中

在CUDA编程中，如果我们要使用共享内存，我们需要将全局内存中的数据带入共享内存。线程用于传输这些数据。复制到cuda中的共享内存中

我在某处阅读（在线资源），最好不要让块中的所有线程都将数据从全局内存复制到共享内存。这样的想法是有道理的，所有的线程不一起执行。经线中的线程一起执行。但我担心的是所有的经纱不是按顺序执行的。比如说，一个带有线程的块被分成3个warp：war p0（0-31 threads），warp 1（32-63 threads），warp 2（64-95 threads）。不能保证warp 0会被首先执行（对不对？）。

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

来源

2013-03-18 user2026934

当块被分配给多处理器时，它驻留在那里，直到该块中的所有线程都完成，并且在此期间，warp调度程序在具有ready操作数的warp之间混合。因此，如果多处理器上有一个块有三个warp，而且只有一个warp从全局到共享内存获取数据，而另外两个warps保持空闲状态并且可能在等待012B的屏障，则不会发生任何事情，并且仅受全局内存延迟的限制无论如何，你会如何。提取完成后，经纱可以继续工作。

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

来源

2013-03-18 00:54:39 stuhlo

要使用单经加载一个共享存储阵列，只是做这样的事情：

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

    // rest of kernel follows 
}

[免责声明：写在浏览器中，从来没有测试，在风险自负]

的关键这里指的是使用__syncthreads()来确保块中的所有线程都等待，直到执行加载到共享内存的warp完成加载。我发布的代码使用了第一个warp，但是您可以通过将warpSize中的线程索引除以块来计算warp数。我也假定了一个一维块，计算二维或三维块中的线索索是微不足道的，所以我把它作为练习给读者。

来源

2013-03-18 07:03:32 talonmies

是warpSize是CUDA内核的已知内部变量吗？ “像threadIDx，blockIdx.etc ...” 我知道它是32，但从变量中获取它对未来的GPU来说更安全:) – 2013-11-24 07:42:03

复制到cuda中的共享内存中

回答

相关问题