0

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

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

4

1 回答 1

3

TL；DR：是的。使用下面的功能。

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

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

但是，如果 CUDA 编译器对我们“隐藏”了寄存器使用，我们怎么能躲在幕后，真正坚持使用它们，用那些%-prefixed 名称访问它们呢？好吧，方法如下：

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

于 2017-02-17T23:52:54.557 回答