用CUDA写出比Numpy更快的规约求和函数
目录 技术背景 CUDA的atomic运算 总结概要 版权声明 技术背景
在前面的几篇博客中我们介绍了在Python中使用Numba来写CUDA程序的一些基本操作和方法,并且展示了GPU加速的实际效果。在可并行化的算法中,比如计算两个矢量的加和,或者是在分子动力学模拟领域中的查找近邻表等等,都是可以直接并行的算法,而且实现起来难度不大。而有一种情况是,如果我们要计算的内容的线程之间互相存在依赖,比方说最常见的,计算一个矩阵所有元素的和。CUDA的atomic运算
正如前面所提到的问题,如何去计算一个矩阵所有元素之和呢?具体问题可以表述为:
S= i,jAi,j
对于此类的问题,如果我们像普通的CUDA并行操作一样,直接创建一个S变量,然后直接在线程和分块上直接把每一个矩阵元素加到这个S变量中,那么会出现一种情况:在线程同步时,存在冲突的线程是无法同时加和成功的,也就是说,这种情况下虽然程序不会报错,但是得到的结果是完全错误的。对于此类情况,CUDA官方给出了atomic运算这样的方案,可以保障线程之间不被干扰:import numpy as np from numba import cuda from numba import vectorize cuda.select_device(1) @cuda.jit def ReducedSum(arr, result): i, j = cuda.grid(2) cuda.atomic.add(result, 0, arr[i][j]) if __name__ == "__main__": import time np.random.seed(2) data_length = 2**10 arr = np.random.random((data_length,data_length)).astype(np.float32) print (arr) arr_cuda = cuda.to_device(arr) np_time = 0.0 nb_time = 0.0 for i in range(100): res = np.array([0],dtype=np.float32) res_cuda = cuda.to_device(res) time0 = time.time() ReducedSum[(data_length,data_length),(1,1)](arr_cuda,res_cuda) time1 = time.time() res = res_cuda.copy_to_host()[0] time2 = time.time() np_res = np.sum(arr) time3 = time.time() if i == 0: print ("The error rate is: ", abs(np_res-res)/res) continue np_time += time3 - time2 nb_time += time1 - time0 print ("The time cost of numpy is: {}s".format(np_time)) print ("The time cost of numba is: {}s".format(nb_time))
这里需要重点关注的就是用CUDA实现的简单函数ReducedSum ,这个函数中调用了CUDA的atomic.add 方法,用这个方法直接替代系统内置的加法,就完成了所有的操作。我们将这个函数的运行时间去跟np.sum 函数做一个对比,结果如下:$ python3 cuda_reduced_sum.py [[0.4359949 0.02592623 0.5496625 ... 0.3810055 0.6834749 0.5225032 ] [0.62763107 0.3184925 0.5822277 ... 0.89322233 0.7845663 0.4595605 ] [0.9666947 0.16615923 0.6931703 ... 0.29497907 0.63724256 0.06265242] ... [0.96224505 0.36741972 0.6673239 ... 0.3115176 0.7561843 0.9396167 ] [0.781736 0.28829736 0.38047555 ... 0.15837361 0.00392629 0.6236886 ] [0.03247315 0.3664344 0.00369871 ... 0.0205253 0.15924706 0.8655231 ]] The error rate is: 4.177044e-06 The time cost of numpy is: 0.027491092681884766s The time cost of numba is: 0.01042938232421875s
在GPU的计算中,会有一定的精度损失,比如这里的误差率就在1e-06级别,但是运行的速度要比numpy的实现快上2倍!总结概要
我们知道GPU加速在可并行化程度比较高的算法中,能够发挥出比较大的作用,展示出明显的加速效果,而对于一些线程之间存在依赖这样的场景就不一定能够起到很大的加速作用。CUDA官方针对此类问题,提供了atomic的内置函数解决方案,包含有求和、求最大值等常用函数。而这些函数的特点就在于,线程与线程之间需要有一个时序的依赖关系。就比如说求最大值的函数,它会涉及到不同线程之间的轮询。经过测试,CUDA的这种atomic的方案,实现起来非常方便,性能也很乐观,相比于自己动手实现一个不断切割、递归的规约函数,还是要容易快捷的多。版权声明
本文首发链接为:https://www.cnblogs.com/dechinphy/p/gpu-sum.html
作者ID:DechinPhy
更多原著文章请参考:https://www.cnblogs.com/dechinphy/
打赏专用链接:https://www.cnblogs.com/dechinphy/gallery/image/379634.html
腾讯云专栏同步:https://cloud.tencent.com/developer/column/91958
手机成了掌上明珠,小孩反而被冷漠了现在除了工作时间放下手机工作外,不管在路上,在车上,还是在家里,哪怕就是在旅游在玩,都是机不离手,手机真正的成为掌上明珠。手机在手,时光总是感觉莫名的觉得过的快,可能9点拿起手机,
浅谈区块链游戏必经的四大发展阶段(三)第二部分,我介绍一下区块链游戏中的Token经济,大家在行业里都在谈,区块链游戏里面虽然说区块链游戏整个体量可能比起整个互联网来说很小,但是游戏是一个天然生态,毋庸置疑区块链将改变
区块链游戏中,通证经济会带来哪些革新?虽然不是所有区块链项目都必须有通证,但没有通证经济作为基础的项目,本质上和中心化项目没什么区别。于是如果你仔细研究各大区块链游戏中的白皮书,它们不约而同一定都会有一定篇幅详细描绘游
2021年了,跨境电商还能做吗?2021年跨境电商还能做吗?在谈论这个问题之前,我觉得还是要先看一下当前的社会环境。许多人都觉得2021年是极其不平凡的一年,疫情复燃洪水四起。特别是我所在的城市郑州,整个7月份和
我为什么坚持留在郑州2019年已经过去了一半,我从没有像现在这样感觉时间快的不像话,大概是到了一定的年龄,都会觉得时间不够用吧。究其根本,是在有限的时间里,还有太多的事情没有完成。四年前我来到郑州的时
华为Mate40Pro系列鸿蒙版本迎来更新了,是一个系统重要补丁包,大小为31。29MB,更新后系统版本依旧为138。更新日志亲爱的用户,本次更新优化了部分场景系统功耗和稳定性,推荐您进行更新。更新注意事项1。本次更新不会删
华为Mate20Pro系列鸿蒙版本迎来更新了华为Mate20Pro系列鸿蒙系统版本也迎来更新了,是一个系统补丁包,大小为30。08MB,更新后系统版本依旧为136。更新日志亲爱的用户,本次更新优化了部分场景系统的功耗和稳定性
华为Mate30Pro5G系列鸿蒙版本迎来更新了华为Mate30Pro5G系列鸿蒙版本迎来更新了,是比较大的更新,大小有1。24GB,更新内容也比较多,更新后系统版本为166。更新包1更新日志亲爱的用户,本次更新相机将新增多机位
家用中央空调漏水怎么办,如何处理?中央空调新装的不久在运行之后,发现有些地方开始渗水或漏水,或有些地方用过一段时间也开始渗水,导致天花板吊顶楼板等地方漏水,这是什么原因呢?新装的中央空调开始漏水,这很可能是在安装的
中央空调电流过大的原因中央空调在运行中有额定的标称值,启动电流一般是运行电流的25倍,不管是运行电流还是启动电流都不能超过标称电流的10,如果超过了就一定有故障的产生,或者不正常的现象,这时要及时停机检
开利中央空调压缩机烧毁怎么办?如何处理?开利中央空调为制冷行业中的知名品牌,全球占有量数年来一直都是遥遥领先,品质和使用体验是有口皆碑。开利中央空调的压缩机分为开利火塞压缩机开利螺杆压缩机开利离心压缩机。占领大中小型的建