我正在尝试使用CUDA开发一个小程序,但由于它是SLOW,我做了一些测试并用Google搜索了一下.我发现虽然单个变量默认存储在本地线程内存中,但数组通常不存在.我想这就是为什么它花了这么多时间来执行.现在我想知道:因为本地线程内存至少应该是16KB,因为我的数组就像52个字符长,有没有办法(语法请:))将它们存储在本地内存中？
不应该是这样的:
__global__ my_kernel(int a)
{
__local__ unsigned char p[50];
}
我的内核调用因"内存不足"而失败.它大量使用了堆栈框架,我想知道这是否是它失败的原因.
使用--ptxas-options = -v调用nvcc时,它会打印以下配置文件信息:
150352 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 59 registers, 40 bytes cmem[0]
硬件:GTX480,sm20,1.5GB设备内存,48KB共享内存/多处理器.
我的问题是分配的堆栈帧在哪里:在共享,全局内存,常量内存,......？
我尝试了每个块1个线程,以及每个块32个线程.同样"内存不足".
另一个问题:如果寄存器的总数不超过多处理器上的可用寄存器数(我的卡为32k),则只能扩大驻留在一个多处理器上的线程数.类似的东西是否适用于堆栈帧大小？
通常建议降低每个线程的寄存器压力以增加 warp 占用，从而提供更大的机会通过 warp 级多线程 (TLP) 隐藏延迟。为了减少寄存器压力，可以使用更多的每个线程本地内存或每个线程块共享内存。CUDA nvcc 编译器也可以强制每个线程使用更少的寄存器。这种方法对于具有良好算术延迟的工作负载很有用，即 ALU 操作与内存 r/w 访问请求的比率很高。然而，对于计算量非常少且内存访问更频繁的延迟关键应用程序，这种方法实际上往往会降低性能。
在此类延迟关键应用程序的情况下，将尽可能多的数据放入片上寄存器或共享内存中更有意义，然后在用全局内存中的下一个数据块替换它之前尽可能多地使用它。当然，通过增加寄存器压力，warp 占用率降低，但现在我们使用快速片上寄存器隐藏片外存储器延迟。增加每个线程寄存器使用的方法是通过展开循环或计算每个线程更多的输出数据来增加 ILP（这也基本上通过对更多输入执行相同的工作来增加 ILP）。这种方法基本上是由 Volkov（Better Performance at Lower Occupancy）提出的。
现在 nvcc 编译器驱动程序有一个名为 maxrregcount 的命令行选项，它允许更改每个线程寄存器的使用。使用此选项一次可以强制编译器减少每个线程寄存器的使用，但不能强制它增加它。我有一个案例，我想增加每个线程寄存器的使用，但我无法在内核中展开循环，因为循环边界是数据相关且动态的。到目前为止，我已经尝试了一些技巧，但我已经没有关于如何增加每个线程寄存器使用的想法。任何人都可以建议增加单个 CUDA 线程的寄存器使用率的方法吗？