如果我使用共享内存，可以分配多少块？
How many blocks can be allocated if i use shared memory?
我是 Cuda 编程新手。
我可以访问设备 "Tesla K10"。
我正在处理一个 复杂问题 ，每个 问题 实例需要大约 20 KB 的内存。
现在，由于 cuda 提供并行化，我决定每个块使用 96 个线程（记住扭曲）来解决问题的一个实例。
现在的问题是我有非常多的 问题 需要解决（比如说超过 1,600,000）。
我知道即使在全局内存中也无法满足如此大的内存需求（在我的例子中是 3.5 GBytes，如下面的 DeviceQuery 输出所示）所以我必须一次解决几个问题。
此外，我已经将每个问题映射到每个块以解决问题的实例。
目前我可以解决大量全局内存中数据的问题。但是共享内存比全局内存快，所以我计划使用 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
...
首先，快速总结一下我是否理解正确：
- 你有大约 150 万个问题要解决，这些是完全独立的，即令人尴尬的并行
- 每个问题都有约 20 KB 的数据集
解决这整个问题需要超过 30 GB 的内存，因此很明显您需要将问题集分成几批。使用您的 4 GB 卡（约 3.5 GB 可用于 ECC 等），您可以随时解决大约 150,000 个问题。如果您要对这些进行双缓冲以允许下一批与当前批的计算同时传输，那么您将在一个批中查看 75K 个问题（如果您需要 space 输出等，则可能更少） .
首先要考虑的重要事情是您是否可以并行处理每个问题，即是否可以将多个线程分配给单个问题？如果是这样，那么您应该考虑分配一个线程块来解决单个问题，使用共享内存可能值得考虑，尽管您会将每个 SM 的占用限制为两个块，这可能会损害性能。
如果您不能在一个问题中并行化，那么您不应该考虑共享内存，因为正如您所说，您会将自己限制为每个 SM 两个线程（从根本上消除了 GPU 计算的优势）。相反，您需要确保全局内存中的数据布局能够实现合并访问 - 这很可能意味着使用 SoA（数组结构）布局而不是 AoS（结构数组）。
你的第二个问题有点令人困惑，不清楚你是指 "block" 在 GPU 上下文中还是在问题上下文中。但是从根本上说，如果您在内核代码中声明了一个 20 KB 的 __shared__ 数组，那么该数组将为每个块分配一次，并且每个块都将具有相同的基址。
根据 OP 的评论更新
GPU包含多个SM，每个SM都有一小块物理内存，既用于L1，也用于共享内存。在您的情况下，K10，每个 SM 有 48 KB 可用作共享内存，这意味着 SM 上所有 执行 的块随时可以在它们之间使用最多 48 KB。由于每个块需要 20 KB，因此您可以随时在 SM 上执行最多两个块。这不会影响您可以在启动配置中设置多少块，它只会影响它们的调度方式。 This answer talks in a bit more detail (albeit for a device with 16 KB per SM) and this (very old) answer explains a little more, although probably the most helpful (and up-to-date) info is on the CUDA education pages.
我是 Cuda 编程新手。 我可以访问设备 "Tesla K10"。 我正在处理一个 复杂问题 ，每个 问题 实例需要大约 20 KB 的内存。 现在，由于 cuda 提供并行化，我决定每个块使用 96 个线程（记住扭曲）来解决问题的一个实例。 现在的问题是我有非常多的 问题 需要解决（比如说超过 1,600,000）。 我知道即使在全局内存中也无法满足如此大的内存需求（在我的例子中是 3.5 GBytes，如下面的 DeviceQuery 输出所示）所以我必须一次解决几个问题。
此外，我已经将每个问题映射到每个块以解决问题的实例。
目前我可以解决大量全局内存中数据的问题。但是共享内存比全局内存快，所以我计划使用 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
...
首先，快速总结一下我是否理解正确：
- 你有大约 150 万个问题要解决，这些是完全独立的，即令人尴尬的并行
- 每个问题都有约 20 KB 的数据集
解决这整个问题需要超过 30 GB 的内存，因此很明显您需要将问题集分成几批。使用您的 4 GB 卡（约 3.5 GB 可用于 ECC 等），您可以随时解决大约 150,000 个问题。如果您要对这些进行双缓冲以允许下一批与当前批的计算同时传输，那么您将在一个批中查看 75K 个问题（如果您需要 space 输出等，则可能更少） .
首先要考虑的重要事情是您是否可以并行处理每个问题，即是否可以将多个线程分配给单个问题？如果是这样，那么您应该考虑分配一个线程块来解决单个问题，使用共享内存可能值得考虑，尽管您会将每个 SM 的占用限制为两个块，这可能会损害性能。
如果您不能在一个问题中并行化，那么您不应该考虑共享内存，因为正如您所说，您会将自己限制为每个 SM 两个线程（从根本上消除了 GPU 计算的优势）。相反，您需要确保全局内存中的数据布局能够实现合并访问 - 这很可能意味着使用 SoA（数组结构）布局而不是 AoS（结构数组）。
你的第二个问题有点令人困惑，不清楚你是指 "block" 在 GPU 上下文中还是在问题上下文中。但是从根本上说，如果您在内核代码中声明了一个 20 KB 的 __shared__ 数组，那么该数组将为每个块分配一次，并且每个块都将具有相同的基址。
根据 OP 的评论更新
GPU包含多个SM，每个SM都有一小块物理内存，既用于L1，也用于共享内存。在您的情况下，K10，每个 SM 有 48 KB 可用作共享内存，这意味着 SM 上所有 执行 的块随时可以在它们之间使用最多 48 KB。由于每个块需要 20 KB，因此您可以随时在 SM 上执行最多两个块。这不会影响您可以在启动配置中设置多少块，它只会影响它们的调度方式。 This answer talks in a bit more detail (albeit for a device with 16 KB per SM) and this (very old) answer explains a little more, although probably the most helpful (and up-to-date) info is on the CUDA education pages.