在内核中声明共享内存变量答案

【问题标题】：declaring shared memory variables in a kernel在内核中声明共享内存变量
【发布时间】：2013-01-30 02:50:45
【问题描述】：

我有一个关于共享变量如何工作的问题。

当我像这样在内核中声明一个共享变量时 __shared__ int array1[N] 每个活动块的每个唯一共享内存现在都有一个大小为 N 的 array1 实例。这意味着每个活动块的每个共享内存现在都分配了 N*sizeof(int) 个字节。对于计算能力为 1.3 的 gpu，N*sizeof(int) 最多必须为 16KB。

所以，假设以上是正确的，并使用像这样在主机上分配的 2D 线程和 2D 块：

dim3 block_size(22,22);
dim3 grid_size(25,25);

我将有 25x25 个 array1 实例，每个实例的大小为 N*sizeof(int)，并且可以访问块的每个共享内存的最多线程是 22x22。这是我最初的问题，并得到了回答。

问：当我给array1赋值时

array1[0]=1;

那么所有活动块是否立即在它们自己的共享内存中分配该值？

【问题讨论】：

你的假设是正确的。但问题是什么？
既然我确认了我的假设，我想问一下，当我使用像这样 dim3 block_size(22,22); 在主机上声明的 2D 线程和 2D 块时dim3 grid_size(25,25);我有多少 array1 实例。 25 还是 25*25？有多少线程可以访问它们中的每一个？谢谢
将有 (25 x 25) 个块，因此每个块共有 625 个单独的共享内存实例。块的共享内存只能由该块的线程访问。所以 (22 x 22) = 484 个线程将能够访问每个共享内存实例。
@sgar91：从来没有那么多独立的静态共享内存分配实例。每个 SM 的并发块数乘以运行内核的 SM 数。
如果调整 L1/shared 拆分，N*sizeof(int) 在 Fermi 和 Kepler 上可以大于 16KB。

标签： cuda shared-memory

【解决方案1】：

每个块将始终分配自己的共享内存数组。因此，如果您启动 25x25 块，您最终将在共享内存中创建 25x25 数组。

然而，这并不意味着所有这些数组都会同时存在，因为不能保证所有块同时存在。活动块的数量取决于运行它的 GPU 的实际型号。 GPU 驱动程序将尝试启动尽可能多的块，并且额外的块将在之前的块结束它们的工作后运行。

N*sizeof(int) 的最大值取决于卡的计算能力和 L1 缓存配置。它可以在：8KB、16KB、32KB 和 48KB 之间变化。

回答您的最后一个问题 - 每个共享数组都对属于相应块的所有线程可见。在您的情况下，相应的 22x22 线程将可以看到每个共享数组。

【讨论】：