我的内核代码可以告诉它有多少共享内存可用吗？答案

【问题标题】：Can my kernel code tell how much shared memory it has available?我的内核代码可以告诉它有多少共享内存可用吗？
【发布时间】：2017-02-17 23:52:54
【问题描述】：

运行设备端 CUDA 代码是否有可能知道为运行内核网格的每个块分配了多少（静态和/或动态）共享内存？

在主机端，您知道启动的内核拥有（或将拥有）多少共享内存，因为您自己设置了该值；但是设备端呢？在该大小的上限内很容易编译，但该信息不可用（除非明确传递）给设备。是否有获取它的 on-GPU 机制？ CUDA C Programming Guide 似乎没有讨论这个问题（在共享内存部分内外）。

【问题讨论】：

标签： cuda gpgpu ptx gpu-shared-memory

【解决方案1】：

TL;DR：是的。使用下面的函数。

有可能：内核代码可以在特殊寄存器中获得该信息：%dynamic_smem_size 和 %total_smem_size。

通常，当我们编写内核代码时，我们不需要知道特定的寄存器（特殊或其他）——我们编写 C/C++ 代码。即使我们确实使用了这些寄存器，CUDA 编译器也会通过保存它们的值的函数或结构对我们隐藏它。例如，当我们使用值threadIdx.x 时，我们实际上是在访问特殊寄存器%tid.x，它为块中的每个线程设置不同。当您查看已编译的 PTX 代码时，您可以看到这些寄存器“正在运行”。 ArrayFire 写了一篇不错的博客文章，其中包含一些工作示例：Demystifying PTX code。

但是如果 CUDA 编译器对我们“隐藏”寄存器使用，我们怎么能躲在幕后，坚持使用它们，使用带有 % 前缀的名称来访问它们？好吧，方法如下：

__forceinline__ __device__ unsigned dynamic_smem_size()
{
    unsigned ret; 
    asm volatile ("mov.u32 %0, %dynamic_smem_size;" : "=r"(ret));
    return ret;
}

和%total_smem_size 的类似功能。该功能使编译器添加显式的 PTX 指令，就像asm 可用于主机代码直接发出 CPU 汇编指令一样。这个函数应该总是内联的，所以当你分配时

x = dynamic_smem_size();

您实际上只是将特殊寄存器的值分配给x。

【讨论】：