【问题标题】:How to allocate all of the available shared memory to a single block in CUDA?如何将所有可用的共享内存分配给 CUDA 中的单个块?
【发布时间】:2013-01-06 03:42:42
【问题描述】:

我想将 SM 的所有可用共享内存分配给一个块。我这样做是因为我不希望将多个块分配给同一个 SM。 我的 GPU 卡有 64KB(共享+L1)内存。在我当前的配置中,48KB 分配给共享内存,16KB 分配给 L1。 我编写了以下代码来使用所有可用的共享内存。

   __global__ void foo()
{

  __shared__ char array[49152];
...

}

我有两个问题:

  1. 如何确保所有共享内存空间都已用完?
  2. 我可以将“48K”增加到一个更高的值(不会出现任何错误或警告)。有没有人可以证明这一点?

提前致谢,

伊曼

【问题讨论】:

  • 当我用一个太大的(静态分配的)共享内存大小进行编译时,我得到一个编译错误(来自 ptxas)。请提供共享内存分配过大的代码的完整可编译示例,以及您用于编译的命令行,以进一步调查您的问题 #2。如果您按照 Eugene 的建议在内核启动时切换到动态分配共享内存,那么如果请求的共享内存太大,您应该会收到运行时错误(您是否正在检查错误?)。

标签: memory cuda gpu


【解决方案1】:
  1. 您可以从cudaDeviceProp::sharedMemPerBlock 读取可用设备共享内存的大小,您可以通过调用cudaGetDeviceProperties 获得
  2. 您不必指定数组的大小。相反,您可以动态传递共享内存的大小作为第三个内核启动参数。

“时钟”CUDA SDK 示例说明了如何在启动时指定共享内存大小。

【讨论】:

  • 如果您的应用程序不支持并发内核执行,那么您只需分配 1/2 * MAX_SHARED_MEMORY_PER_SM + 1 字节的共享内存即可将占用限制为每个 SM 1 个块。
猜你喜欢
  • 1970-01-01
  • 1970-01-01
  • 2013-05-26
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
  • 1970-01-01
相关资源
最近更新 更多