CUDA 块级共享寄存器答案

【问题标题】：CUDA Block-level Shared RegistersCUDA 块级共享寄存器
【发布时间】：2021-08-12 05:00:23
【问题描述】：

有什么方法可以让我使用 registers 以便线程束/块可以访问，就像 shared memory 一样？

另一个假设是我已经有了每个扭曲/块的大小，例如，我需要64xsizeof(float) 寄存器。

主要原因是我发现在应用程序需要非合并访问的情况下，从 warp 到共享内存的随机访问非常慢。

【问题讨论】：

标签： cuda gpu nvidia

【解决方案1】：

一般不会。

您没有强制或指示编译器将项目放入寄存器的机制。例如，C++ register 关键字不这样做。寄存器的使用由编译器控制。
即使可以，寄存器也不会像共享内存那样“共享”，也不会像 warp/threadblock 中的多个线程可见。寄存器一旦分配给拥有它们的线程，就仅在该线程可见。它们不能对经线或线程块中的任何其他线程可见。

如果您只是需要每个线程的本地存储，而不是在线程之间共享，并且您的使用满足了许多其他要求，那么您可以“哄骗”编译器将这些项目放入寄存器中。 This forum thread 提供了一个示例，说明它是如何完成的，以及它的外观。同样，一旦编译器清楚地看到这样做的方式，并确定有充分的理由这样做（即“相信”可能的性能改进），编译器将自动将项目放入寄存器中。您无法直接控制它。

【讨论】：

是否可以在我的 CUDA 代码中嵌入一些 PTX 程序集？例如，在每个经线的开始。为每个块/扭曲分配一组寄存器；每个线程将根据它们的线程laneid使用这些寄存器（即，每个线程将知道它将读取和写入哪个寄存器）。对此有何建议或建议？
当然可以embed PTX assembly in CUDA code。在cuda 标签上有很多这样做的例子。但这并没有改变这个答案。 PTX 不是强制编译器以某种方式使用寄存器的东西。实际寄存器使用情况由ptxas 工具确定，该工具采用 PTX 并将其转换为 SASS，即在 GPU 上运行的实际代码。 ptxas 是一个优化编译器，而不仅仅是一个汇编器。此外，PTX 不允许在线程之间共享寄存器。
好吧，我明白了，我想知道 wmma::fragment 中的 Tensor Core 计算的情况，其中来自线程束的所有寄存器都构建了这个片段，我怎么知道我从每个寄存器访问哪个寄存器线？例如，在线程中，当我调用a_frag.x[t].
没有办法先验地知道将使用哪些寄存器。（当然，您可以使用二进制实用程序检查生成的 SASS。）此外，加载片段不会改变我提到的任何事情。寄存器不会在 warp 中的线程之间共享。片段实际上是 warp 中每个线程的一组寄存器。经线中的任何线程都无法访问整个片段（一旦加载）。它只能访问它自己的片段部分。
所以没有办法知道a_frag.x[t]的映射 --> 片段中元素的索引？