c++ - 如果我使用共享内存，可以分配多少 block ？

Question

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

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

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

1) 现在我的困惑是这将只允许我一次将 2 个问题加载到共享内存中以解决(即 40KBytes < 48KBytes 共享内存)。 (基于我对cuda的理解，如有不妥请指正)。

2) 如果我在内核中用这 20 KBytes 声明数组是否意味着这 (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 个问题(如果您需要输出空间等，可能会更少)。

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

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

您的第二个问题有点令人困惑，不清楚您指的是 GPU 上下文中的“ block ”还是问题上下文中的“ block ”。然而，如果您在内核代码中声明一个 20 KB 的 __shared__ 数组，那么该数组将在每个 block 中分配一次，并且每个 block 都将具有相同的基地址。

根据 OP 的评论更新

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