31

我正在尝试用 CUDA 开发一个小程序，但由于它很慢，我做了一些测试并用谷歌搜索了一下。我发现虽然单个变量默认存储在本地线程内存中，但数组通常不是。我想这就是为什么执行需要这么多时间的原因。现在我想知道：由于本地线程内存应该至少为 16KB，而且我的数组只有 52 个字符长，有没有办法（请使用语法 :)）将它们存储在本地内存中？

不应该是这样的：

__global__ my_kernel(int a)
{
  __local__ unsigned char p[50];
}
4

5 回答 5

82

数组、本地内存和寄存器

这里对“本地内存”的定义存在误解。CUDA 中的“本地内存”实际上是具有交错寻址的全局内存（实际上应该称为“线程本地全局内存”）（这使得并行迭代数组比将每个线程的数据阻塞在一起要快一点）。如果你想让事情变得非常快，你可以使用共享内存，或者更好的是寄存器（尤其是在最新的设备上，每个线程最多可以有 255 个寄存器）。解释整个CUDA 内存层次结构超出了本文的范围。让我们专注于快速进行小数组计算。

小数组，就像变量一样可以完全存储在寄存器中。然而，在当前的 NVIDIA 硬件上，将数组放入寄存器是很困难的。为什么？因为寄存器需要非常小心的处理。如果您没有完全正确地执行此操作，您的数据将最终存储在本地内存中（同样，这实际上是全局内存，这是您拥有的最慢的内存）。CUDA 编程指南第5.3.2 节告诉您何时使用本地内存：

本地内存

本地内存访问仅发生在变量类型限定符中提到的一些自动变量。编译器可能放置在本地内存中的自动变量是：

  1. 无法确定它们是否以常数索引的数组，
  2. 会占用过多寄存器空间的大型结构或数组，
  3. 如果内核使用的寄存器多于可用寄存器（这也称为寄存器溢出），则为任何变量。

寄存器分配是如何工作的？

请注意，寄存器分配是一个极其复杂的过程，这就是为什么您不能（也不应该）干预它的原因。相反，编译器会将 CUDA 代码转换为 PTX 代码（一种字节码），它假定机器具有无限多的寄存器。您可以编写内联 PTX，但它不会对注册分配做太多。PTX 代码是与设备无关的代码，它只是第一阶段。在第二阶段，PTX 将被编译成设备汇编代码，称为 SASS。SASS 代码具有实际的寄存器分配。SASS 编译器和它的优化器也将是决定变量是在寄存器中还是在本地内存中的最终权威。您所能做的就是尝试了解 SASS 编译器在某些情况下的作用，并将其用于您的优势。Nsight 中的代码关联视图可以帮助您解决这个问题（见下文）。

寄存器不足

附录 G，第 1 节告诉您一个线程可以拥有多少个寄存器。查找“每个线程的最大 32 位寄存器数”。为了解释该表，您必须了解您的计算能力（见下文）。不要忘记寄存器用于各种事物，并且不仅仅与单个变量相关。最高 CC 3.5 的所有设备上的寄存器均为 32 位。如果编译器足够聪明（并且 CUDA 编译器不断变化），它可以例如将多个字节打包到同一个寄存器中。Nsight 代码相关视图（请参阅下面的“分析内存访问”）也揭示了这一点。

常量与动态索引

虽然空间限制是注册数组的明显障碍，但很容易监督的事实是，在当前硬件（Compute Capability 3.x 及更低版本）上，编译器将任何数组放置在本地内存中，通过动态索引。动态索引是编译器无法计算的索引。使用动态索引访问的数组不能放在寄存器中，因为寄存器必须由编译器确定，因此实际使用的寄存器不能依赖于运行时确定的值。例如，给定一个数组arr，当且仅当它是一个常量，或者仅依赖于常量时，arr[k]它才是常量索引。k如果k以任何方式依赖于某个非常量值，则编译器无法计算k你得到了动态索引。k在以（小）常数开始和结束的循环中，编译器（很可能）可以展开您的循环，并且仍然可以实现常数索引。

例子

例如，可以在寄存器中对小数组进行排序，但您必须使用排序网络或类似的“硬连线”方法，并且不能只使用标准算法，因为大多数算法都使用动态索引。

很有可能，在下面的代码示例中，编译器将整个aBytes数组保存在寄存器中，因为它不是太大，并且循环可以完全展开（因为循环在恒定范围内迭代）。编译器（很可能）知道每一步都在访问哪个寄存器，因此可以将其完全保存在寄存器中。请记住，没有任何保证。您可以做的最好的事情是使用 CUDA 开发人员工具逐个验证它，如下所述。

__global__
void
testSortingNetwork4(const char * aInput, char * aResult)
{
    const int NBytes = 4;

    char aBytes[NBytes];

    // copy input to local array
    for (int i = 0; i < NBytes; ++i)
    {
        aBytes[i] = aInput[i];
    }

    // sort using sorting network
    CompareAndSwap(aBytes, 0, 2); CompareAndSwap(aBytes, 1, 3); 
    CompareAndSwap(aBytes, 0, 1); CompareAndSwap(aBytes, 2, 3); 
    CompareAndSwap(aBytes, 1, 2); 


    // copy back to result array
    for (int i = 0; i < NBytes; ++i)
    {
        aResult[i] = aBytes[i];
    }
}

分析内存访问

完成后，您通常要验证数据是否实际存储在寄存器中或是否进入本地内存。您可以做的第一件事是告诉您的编译器使用标志为您提供内存统计信息--ptxas-options=-v。分析内存访问的更详细方法是使用Nsight。

Nsight 有许多很酷的功能。Nsight for Visual Studio有一个内置的分析器和一个 CUDA <-> SASS 代码相关视图。该功能在此处进行了说明。请注意，不同 IDE 的 Nsight 版本可能是独立开发的，因此它们的功能可能因不同的实现而异。

如果您按照上面链接中的说明进行操作（请确保在编译时添加相应的标志！），您可以在下方菜单的最底部找到“CUDA Memory Transactions”按钮。在该视图中，您希望发现没有来自仅在相应数组上工作的行（例如，我的代码示例中的CompareAndSwap行）的内存事务。因为如果它不报告这些行的任何内存访问，您（很可能）能够将整个计算保存在寄存器中，并且可能刚刚获得了数千甚至一万百分比的加速（您可能还想检查实际的速度增益，你可以摆脱这个！）。

弄清楚计算能力

为了弄清楚你有多少个寄存器，你需要知道你的设备的计算能力。获取此类设备信息的标准方法是运行设备查询示例。对于 Windows 64 位上的 CUDA 5.5，默认情况下位于C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\Bin\win64\Release\deviceQuery.exe（在 Windows 上，控制台窗口将立即关闭，您可能需要首先打开cmd并从那里运行它）。它在 Linux 和 MAC 上的位置相似。

如果您有 Nsight for Visual Studio，只需转到 Nsight -> Windows -> 系统信息。

不要过早优化

我今天分享这个是因为我最近遇到了这个问题。然而，正如在这个线程中提到的，强制数据进入寄存器绝对不是你想要采取的第一步。首先，确保您真正了解发生了什么，然后逐步解决问题。查看汇编代码当然是一个很好的步骤，但它通常不应该是你的第一步。如果您是 CUDA 新手，CUDA 最佳实践指南将帮助您了解其中的一些步骤。

于 2013-09-27T01:49:36.450 回答
11

你只需要这样：

__global__ my_kernel(int a)
{
    unsigned char p[50];
    ........
}

如果需要，编译器会自动将其溢出到线程本地内存。但请注意，本地内存存储在 GPU 之外的 SDRAM 中，它与全局内存一样慢。因此，如果您希望这会带来性能改进，那么您可能会感到失望......

于 2012-04-24T11:47:07.997 回答
1

〜对于将来遇到这种情况的人〜

简而言之，要为每个线程创建一个数组，您需要在设备内存中创建它们。为此，可以为每个线程分配一点共享内存。必须特别注意防止冲突或性能下降。

以下是Maxim Milakov 在 2015 年发表的一篇英伟达博客文章中的一个示例：

// Should be multiple of 32
#define THREADBLOCK_SIZE 64 
// Could be any number, but the whole array should fit into shared memory 
#define ARRAY_SIZE 32 

__device__ __forceinline__ int no_bank_conflict_index(int thread_id, int logical_index)
{
    return logical_index * THREADBLOCK_SIZE + thread_id;
}

__global__ void kernel5(float * buf, int * index_buf)
{
    // Declare shared memory array A which will hold virtual 
    // private arrays of size ARRAY_SIZE elements for all 
    // THREADBLOCK_SIZE threads of a threadblock
    __shared__ float A[ARRAY_SIZE * THREADBLOCK_SIZE]; 
    ...
    int index = index_buf[threadIdx.x + blockIdx.x * blockDim.x];

    // Here we assume thread block is 1D so threadIdx.x 
    // enumerates all threads in the thread block
    float val = A[no_bank_conflict_index(threadIdx.x, index)];
    ...
}
于 2019-07-23T19:15:53.893 回答
-1

您正在混淆本地和注册内存空间。

单个变量和恒定大小的数组自动保存在芯片上的寄存器空间中，几乎没有读写成本。

如果您超过每个多处理器的寄存器数量，它们将存储在本地内存中。

本地内存驻留在全局内存空间中，具有相同的慢速读写操作带宽。

#DEFINE P_SIZE = 50

__global__ void kernel()
{
    unsigned char p[P_SIZE];
}
于 2012-04-24T13:00:22.257 回答
-1

您要查找的关键字是__shared__. 大型数组不适合共享内存空间，但编译器应该将共享内存用于固定大小的小型数组，就像在这种情况下一样。您可以使用__shared__关键字来确保发生这种情况。如果超过块的最大共享内存量，您将看到编译时错误。

于 2012-04-24T13:31:25.107 回答