Mat*_*nti 23 memory arrays cuda gpu-local-memory
我正在尝试使用CUDA开发一个小程序,但由于它是SLOW,我做了一些测试并用Google搜索了一下.我发现虽然单个变量默认存储在本地线程内存中,但数组通常不存在.我想这就是为什么它花了这么多时间来执行.现在我想知道:因为本地线程内存至少应该是16KB,因为我的数组就像52个字符长,有没有办法(语法请:))将它们存储在本地内存中？
不应该是这样的:
__global__ my_kernel(int a)
{
__local__ unsigned char p[50];
}
Dom*_*omi 64
这里有关于"本地记忆"定义的误解.CUDA中的"本地存储器"实际上是全局存储器(实际上应该称为"线程局部全局存储器"),具有交错寻址(这使得并行迭代比使每个线程的数据一起阻塞更快).如果你想让事情变得非常快,你想要使用共享内存,或者更好的是寄存器(特别是在每个线程最多可以获得255个寄存器的最新设备上).解释整个CUDA内存层次结构超出了本文的范围.让我们专注于快速进行小数组计算.
小数组就像变量一样可以完全存储在寄存器中.然而,在当前的NVIDIA硬件上,将阵列放入寄存器是困难的.为什么？因为寄存器需要非常小心的处理 如果你不完全正确,你的数据将最终在本地内存中(这也是真正的全局内存,这是你拥有的最慢的内存)." CUDA编程指南"的第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以(小)常数开始和结束的循环中,编译器(最有可能)可以展开循环,并且仍然可以实现持续索引.
例如,可以在寄存器中对小数组进行排序,但必须使用排序网络或类似的"硬连线"方法,并且不能仅使用标准算法,因为大多数算法都使用动态索引.
具有相当高的概率,在下面的代码示例中,编译器将整个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 64bit上的CUDA 5.5,默认情况下位于C:\ ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\Bin\win64\Release\deviceQuery.exe(在Windows上,控制台窗口将立即关闭,您可能需要首先打开cmd并从那里运行它.它在Linux和MAC上具有类似的位置.
如果你有Visual Studio的Nsight,只需转到Nsight - > Windows - > System Info.
我今天分享这个,因为我最近遇到了这个问题.但是,如此线程中所述,强制数据在寄存器中绝对不是您想要采取的第一步.首先,确保您实际了解正在发生的事情,然后逐步解决问题.查看汇编代码肯定是一个很好的步骤,但它通常不应该是您的第一步.如果您是CUDA的新手,CUDA最佳实践指南将帮助您找出其中的一些步骤.
tal*_*ies 11
你需要的只是这个:
__global__ my_kernel(int a)
{
unsigned char p[50];
........
}
如果需要,编译器将自动将其溢出到线程本地内存.但请注意,本地内存存储在GPU的SDRAM中,并且与全局内存一样慢.因此,如果您希望这会带来性能提升,那么您可能会感到失望.....