如果我使用共享内存可以分配多少块？

Question

我是 Cuda 编程的新手。 我可以访问设备“Tesla K10”。 我正在研究一个复杂问题，每个问题实例需要大约 20 KB 的内存。 现在由于 cuda 提供并行化，我决定每个块使用 96 个线程（记住扭曲）来解决问题的一个实例。 现在的问题是我有大量的此类问题需要解决（比如超过 1,600,000 个）。 我知道即使在全局内存中也无法满足如此大的内存需求（在我的情况下是 3.5 GB，如下面的 DeviceQuery 输出所示），因此我必须一次解决几个问题。

此外，我已将每个问题与每个块映射以解决问题的一个实例。

现在我能够解决全局内存中数据的大量问题。但是共享内存比全局快，所以我计划使用 20 KB 的共享内存（每个问题）。

1) 现在我的困惑是，这一次只能将 2 个问题加载到共享内存中来解决（即 40KBytes

2) 如果我在内核中声明这个 20 KB 的数组，这是否意味着这 (20KBytes * number_of_blocks) 将是共享内存使用？ number_of_blocks是指要解决的问题的数量。 我的启动配置是 问题>(...)

您在这方面的所有帮助都将得到高度认可。 提前谢谢你。

***My partial Device Query***

Device : "Tesla K10.G1.8GB"
  CUDA Driver Version / Runtime Version          6.5 / 5.0
  CUDA Capability Major/Minor version number:    3.0
  Total amount of global memory:                 3584 MBytes (3757637632 bytes)
  ( 8) Multiprocessors, (192) CUDA Cores/MP:     1536 CUDA Cores
  GPU Clock rate:                                745 MHz (0.75 GHz)
  Memory Clock rate:                             524 Mhz
  Memory Bus Width:                              2048-bit
  L2 Cache Size:                                 4204060 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 2046), 65536 layers
  Total amount of constant memory:               65536 bytes
  **Total amount of shared memory per block:       49152 bytes**
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  0
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
...

Answer 1

首先，快速总结一下我是否理解正确：

• 您有大约 150 万个问题要解决，这些问题是完全独立的，即令人尴尬的并行
• 每个问题都有一个约 20 KB 的数据集

解决整个问题需要 >30 GB 的内存，因此很明显您需要将问题集拆分为多个批次。使用您的 4 GB 卡（约 3.5 GB 可使用 ECC 等），您可以随时解决大约 150,000 个问题。如果您要将这些缓存加倍以允许下一批与当前批次的计算同时传输，那么您将在一个批次中查看 75K 问题（如果您需要输出空间等可能会更少）。

首先要考虑的是您是否可以并行化每个问题，即有没有办法将多个线程分配给一个问题？如果是这样，那么您应该考虑分配一个线程块来解决单个问题，使用共享内存可能值得考虑，尽管您将每个 SM 的占用限制为两个块，这可能会损害性能。

如果您无法在问题中进行并行化，那么您不应该考虑共享内存，因为正如您所说，您将自己限制为每个 SM 两个线程（从根本上消除了 GPU 计算的好处）。相反，您需要确保全局内存中的数据布局能够实现合并访问 - 这很可能意味着使用 SoA（数组结构）布局而不是 AoS（结构数组）。

您的第二个问题有点令人困惑，不清楚您的意思是在 GPU 上下文中还是在问题上下文中“阻塞”。但是从根本上说，如果您在内核代码中声明一个 20 KB 的 __shared__ 数组，那么该数组将在每个块中分配一次，并且每个块将具有相同的基地址。

跟随 OP 的 cmets 更新

GPU 包含许多 SM，每个 SM 都有一小块物理内存，用于 L1 和共享内存。在您的情况下，K10，每个 SM 有 48 KB 可用作共享内存，这意味着 SM 上的所有块执行在任何时候都可以在它们之间使用最多 48 KB。由于每个块需要 20 KB，因此您可以随时在 SM 上执行最多两个块。这不会影响您可以在启动配置中设置多少块，它只会影响它们的安排方式。 This answer 讲得更详细一些（尽管是针对每个 SM 16 KB 的设备），this (very old) answer 解释得更多，尽管可能最有用（也是最新）的信息在 CUDA education pages 上。