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

Question

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

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

Answer 1

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

有可能：该信息可用于特殊寄存器中的内核代码：%dynamic_smem_size and %total_smem_size。

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

但是，如果 CUDA 编译器 "hides" 向我们注册使用，我们如何才能回到幕后并真正坚持使用它们，并使用那些 % 前缀的名称访问它们？那么，方法如下：

__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.

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

Can my kernel code tell how much shared memory it has available?

cuda

gpgpu

ptx

gpu-shared-memory

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