明显的CUDA魔法
Apparent CUDA magic
我正在使用CUDA(实际上我正在使用pyCUDA,如果差异重要的话)并对数组执行一些计算。我正在启动一个具有320*600线程网格的内核。在内核内部,我使用
声明两个20000个组件的线性数组:float test[20000]
float test2[20000]
对于这些数组,我执行简单的计算,例如用常数值填充它们。关键是内核正常执行, 正确执行的计算(您可以看到这用test的随机组件填充数组并将该数组从设备发送到主机)。问题是我的NVIDIA卡只有2GB的内存,分配数组test和test2的内存总量是320*600*20000*4字节,远远超过2GB。
这个内存是从哪里来的?CUDA如何在每个线程中执行计算?
感谢您的宝贵时间
本地/堆栈内存需求的实际大小并不像您想象的那样(对于整个网格,一次全部),而是实际上基于@njuffa在这里描述的公式。
基本上,本地/堆栈内存需求的大小是基于您正在运行的设备的最大瞬时容量,而不是网格的大小。
根据njuffa提供的信息,可用的堆栈大小限制(每个线程)是以下选项中较小者:
- 最大本地内存大小(cc2为512KB)。x及以上)
- 可用 GPU内存/(SMs数量)/(每个SMs最大线程数)
对于您的第一个案例:
float test[20000];
float test2[20000];
这个总数是160KB(每个线程),所以我们在每个线程512KB的最大限制之下。第二个极限呢?
GTX 650m有2 cc 3.0 (kepler) SMs(每个kepler SM有192个内核)。因此,上述第二个限制给出,如果所有 GPU内存可用:
2GB/2/2048 = 512KB
(开普勒每个多处理器有2048个最大线程)在这种情况下是相同的极限。但是这假设所有的GPU内存都是可用的。
既然你在评论中建议这个配置失败:
float test[40000];
float test2[40000];
。320KB,我会得出结论,你的实际可用GPU内存是在这个批量分配尝试的点上,高于(160/512)*100%,即高于31%,但低于(320/512)*100%,即低于2GB的62.5%,所以我会得出结论,你的可用GPU内存在这个批量分配请求堆栈帧的时间将小于1.25GB。
你可以试着通过在内核启动之前调用cudaGetMemInfo
来查看是否存在这种情况(尽管我不知道如何在pycuda中这样做)。即使你的GPU开始时是2GB,如果你从它运行显示,你可能从接近1.5GB的数字开始。在内核启动时,在此批量分配请求之前进行的动态(例如cudaMalloc
)和/或静态(例如__device__
)分配都会影响可用内存。
这些都是为了解释一些细节。对你的问题的一般回答是,"魔法"的出现是由于GPU不一定要为网格中的所有线程一次分配堆栈帧和本地内存。它只需要分配设备的最大瞬时容量所需的内容(即SMs *每个SM的最大线程数),这可能是一个明显小于整个网格所需的数字。
- 编译时未启用intel oneApi CUDA支持
- 在cuda线程之间共享大量常量数据
- 为什么即使使用-cudart-static进行编译,库用户仍然需要链接到cuda运行时
- Cuda C++:设备上的Malloc类,并用来自主机的数据填充它
- CUDA内核和数学函数的显式命名空间
- CUDA:统一内存和指针地址的更改
- 调试 CUDA MMU 故障
- 使用 CUDA 和纹理进行图像减法
- 将 2D 推力::d evice_vector 复矩阵传递给 CUDA 内核函数
- 编译 CUDA 与数学函数的叮当
- 为什么 CUDA 不会导致C++代码加速?
- 如何防止 CUDA-GDB 中的<优化输出>值
- 通过Python Distutils(用于Python C扩展)使用可重定位的设备代码编译CUDA代码
- CUDA三角函数中的数学保证
- CUDA 使用共享内存平铺 3D 卷积实现
- CUDA:cudaMallocManage处理退出吗?
- Opencv 加速与 CUDA 在 C++.
- Cuda:具有位集数组的 XOR 单位集
- 用于构建 cuda .so 文件(共享库)的生成文件
- 明显的CUDA魔法