我正在尝试使用 CUDA 开发一个小程序，但由于它很慢，所以我做了一些测试并在 Google 上搜索了一下。我发现虽然单个变量默认存储在本地线程内存中，但数组通常不是。我想这就是为什么执行需要这么多时间的原因。现在我想知道：由于本地线程内存至少应为 16KB，而且我的数组只有 52 个字符长，有没有办法（请使用语法 :)）将它们存储在本地内存中？
不应该是这样的：
__global__ my_kernel(int a)
{
__local__ unsigned char p[50];
}
【问题讨论】：
__constant__ 数组中。详情请参阅编程指南。
标签： arrays memory cuda gpu-local-memory
这里对“本地内存”的定义存在误解。 CUDA 中的“本地内存”实际上是具有交错寻址的全局内存（实际上应该称为“线程本地全局内存”）（这使得并行迭代数组比将每个线程的数据阻塞在一起要快一点）。如果你想让事情变得非常快，你想使用共享内存，或者更好的是，寄存器（特别是在最新的设备上，每个线程最多可以有 255 个寄存器）。解释整个CUDA memory hierarchy 超出了本文的范围。让我们专注于快速进行小型数组计算。
小数组，就像变量可以完全存储在寄存器中一样。然而，在当前的 NVIDIA 硬件上，将数组放入寄存器是很困难的。为什么？因为寄存器需要非常小心的处理。如果您没有完全正确地执行此操作，您的数据将最终存储在本地内存中（同样，这实际上是全局内存，这是您拥有的最慢的内存）。 CUDA Programming Guide, section 5.3.2 告诉您何时使用本地内存：
本地内存
本地内存访问仅发生在变量类型限定符中提到的一些自动变量。编译器可能放置在本地内存中的自动变量是：
- 无法确定是否使用常数索引的数组，
- 会占用过多寄存器空间的大型结构或数组，
- 内核使用的寄存器多于可用寄存器时的任何变量（这也称为寄存器溢出）。
请注意，寄存器分配是一个极其复杂的过程，这就是为什么您不能（也不应该）干预它的原因。相反，编译器会将 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 以（小）常数开始和结束的循环中，编译器（很可能）可以展开您的循环，并且仍然可以实现常数索引。
例如，可以在寄存器中对小数组进行排序，但您必须使用sorting networks 或类似的“硬连线”方法，并且不能只使用标准算法，因为大多数算法都使用动态索引。
很有可能，在下面的代码示例中，编译器将整个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];
}
}
完成后，您通常要验证数据是否实际存储在寄存器中或是否进入本地内存。您可以做的第一件事是tell your compiler to give you memory statistics using the --ptxas-options=-v flag。更详细的分析内存访问的方法是使用Nsight。
Nsight 有许多很酷的功能。 Nsight for Visual Studio 有一个内置的分析器和一个 CUDA SASS 代码相关视图。 The feature is explained here。请注意，不同 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 -> 系统信息。
我今天分享这个是因为我最近遇到了这个问题。然而，正如this thread 中提到的，强制数据进入寄存器绝对不是你想要采取的第一步。首先，确保您真正了解发生了什么，然后逐步解决问题。查看汇编代码当然是一个很好的步骤，但它通常不应该是你的第一步。如果您是 CUDA 新手，CUDA Best Practices Guide 将帮助您了解其中的一些步骤。
【讨论】：
-maxrregcount=72（相对于 94 个寄存器，没有此选项）为 CUDA-raytracer 提供了 1.5 倍的加速。
你只需要这个：
__global__ my_kernel(int a)
{
unsigned char p[50];
........
}
如果需要，编译器会自动将其溢出到线程本地内存。但请注意，本地内存存储在 GPU 之外的 SDRAM 中，它与全局内存一样慢。因此，如果您希望这会带来性能提升，那么您可能会感到失望.....
【讨论】：
~对于以后遇到这个的人~
简而言之，要为每个线程创建一个数组，您需要在设备内存中创建它们。为此，可以为每个线程分配一点 shared 内存。必须特别注意防止冲突或性能下降。
这是一个来自 Maxim Milakov 在 2015 年 nvidia blog post 的示例：
// 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)];
...
}
【讨论】：
您要查找的关键字是 __shared__ 。大型数组不适合共享内存空间，但编译器应该将共享内存用于固定大小的小型数组，就像在这种情况下一样。您可以使用__shared__ 关键字来确保发生这种情况。如果超过一个块的最大共享内存量，您将看到编译时错误。
【讨论】：
__shared__ 内存不仅仅用于共享数据。它可以用作每个线程的暂存器——如果每个线程多次访问相同的数据，它会很有用。它还可以用于有效地重新排序数据，以便可以按顺序完成全局内存访问（有关此示例，请参见转置 CUDA SDK 示例）。
您正在混淆本地和注册内存空间。
单个变量和固定大小的数组自动保存在芯片上的寄存器空间中，几乎没有读写成本。
如果您超过每个多处理器的寄存器数量，它们将被存储在本地内存中。
本地内存驻留在全局内存空间中，读写操作带宽相同。
#DEFINE P_SIZE = 50
__global__ void kernel()
{
unsigned char p[P_SIZE];
}
【讨论】：