c++ - cuda 将大于线程数的数组复制到共享内存

Question

如果内核出于某些其他原因需要 y 线程，那么某人将如何复制大小为 x 的数组。我找到的所有示例都是复制线程数大小的数组，例如:

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;

  s[t] = d[t];
  __syncthreads();

1. 如果我需要我的 s[100000] 并且我需要仅使用 1 个 640 线程 block 来启动我的内核，它会是什么样子？

2. 如果我需要我的 s[100000] 并且我需要用 10 个 64 线程 block 启动我的内核，它会是什么样子？在这里我也很困惑，因为共享内存只用于 1 个 block ，所以我不明白其他 block 的线程如何复制到哪个 block 的共享内存？

Answer 1

How would it look If I need my s[100000] and I need to launch my kernel with just 1 block of 640 threads?

不会的。目前支持的 CUDA 硬件有一个 limit每 block 48kb 或 96kb，因此具有 100000 个元素的整数共享数组是非法的。

但一般来说，如果您有一种设计模式，其中有必要将每个线程的多个数据点加载到共享内存中，那么您会做这样的事情

__global__ void staticReverse(int *d, int n)
{
  __shared__ int s[2048];
  int t = threadIdx.x;

  for(; t < 2048; t += blockDim.x) s[t] = d[t];
  __syncthreads();

因为共享内存是严格的 block 范围，所以您必须使用类似上述设计模式的东西来将数据加载到共享存储。显然，来自其他 block 的线程永远无法访问除它们自己的 block 之外的任何共享内存。