cuda的Shuffle技术以及自定义双精度版本
来源:互联网 发布:华康海报简体mac 编辑:程序博客网 时间:2024/06/08 18:06
还是数组求和问题引起的,发现之前那个版本http://blog.csdn.net/lingerlanlan/article/details/24630511
对于数组的维度是有要求的。因为归约每次变为一半,所以对于线程块的数量和每个线程块线程的数量都要是2的倍数。
今天看到这篇文章https://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/。
对并行归约进行了讨论。目前还没完全读懂,读懂了翻译一下。
现在对刚了解的shuffle技术写一下体会。
这玩意就是使得线程束内的线程可以共享寄存器变量。
比如函数
int __shfl_down(int var, unsigned int delta, int width=warpSize);
有点像在线程间左移变量。
下面用具体例子来说明,
int i = threadIdx.x % 32;int j = __shfl_down(i, 2, 8);这里32指一个线程束的线程数量是32
第一句:
int i = threadIdx.x % 32;
每个线程都有一个变量i,即是线程在所在线程束的id。
第二句:
int j = __shfl_down(i, 2, 8);首先8指明了范围,就是0-7,8-15,16-23,24-31。
2指明了步长。比如i=5的线程,把i值赋值给了i=3的线程中的j变量。本质上就是在一定范围内线程间按照一定的步长来访问另一格线程的寄存器变量。
这幅图很好的说明了
测试例子:
#include <stdio.h>__global__ void kernel(){int i = threadIdx.x % 32;int j = __shfl_down(i, 2, 8);printf("%d:%d\n",i,j);}int main(){kernel<<<1,32>>>();cudaDeviceSynchronize();return 0;}
输出结果:
0:21:32:43:54:65:76:67:78:109:1110:1211:1312:1413:1514:1415:1516:1817:1918:2019:2120:2221:2322:2223:2324:2625:2726:2827:2928:3029:3130:3031:31
注意红色的部分,因为参数8指明了执行范围。
因为库指提供了int和float的shuffle版本,http://docs.nvidia.com/cuda/cuda-c-programming-guide/#warp-shuffle-functions。
双精度的需要自己实现
__device__ inlinedouble __shfl_down(double var, unsigned int srcLane, int width=32) { int2 a = *reinterpret_cast<int2*>(&var); a.x = __shfl_down(a.x, srcLane, width); a.y = __shfl_down(a.y, srcLane, width); return *reinterpret_cast<double*>(&a);}
这个很巧妙的。用两个32位的int来跟64位的double转换。
其实理解这个,关键是要彻底明白计算机存储数据就是若干个0和1。
而这里巧妙的另外一个地方是用到了
reinterpret_cast函数来强制转换。
这让我想起了曾经面试qq后台开发经历,貌似就是实现两个很大整数数的相加,具体多少位忘了,反正超过32位。
应该就是这种思路。
参考资料:
https://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/
- cuda的Shuffle技术以及自定义双精度版本
- cuda版本的word2vec
- cuda的计时技术
- cuda的优化技术
- 双精度浮点型数据运算精度丢失以及数据的格式化问题
- Shuffle阶段:partition分区以及自定义使用注意事项
- mapreduce的执行流程以及shuffle过程
- CUDA之Warp Shuffle详解
- Cordova-iOS自定义插件以及和老版本的差异
- 单精度和双精度的区别
- 单精度与双精度的区别
- 关于单精度 双精度的概念
- 【C++】求取两个整数、三个整数,两个双精度数、三个双精度数的最大值【原创技术】
- js的转换精度问题以及解决方案
- 如何查看显卡支持的CUDA版本
- 历史版本的cuda和cuDNN
- caffe 与 cuda 版本的问题
- 查看cuda和cudnn版本的命令
- 即时定位与地图构建(SLAM)的相关研究
- Jquery-zTree的基本用法-java版本
- SpringMVC 快速入门
- Android软件开发之盘点所有Dialog对话框大合集
- 交换友情链接注意事项
- cuda的Shuffle技术以及自定义双精度版本
- 笑笑十年少
- 全面降低准备金率会引来三大恶果?
- EditBox输入时自动换行
- OCP 1Z0 051 6
- 在Qt中使用opencv库
- Bootstrap风格的按钮 BButton(把图标做成字体,可以省下很多图片资源)
- 看example源码学spark系列(1)-LocalPi
- 电视剧《花千骨》开机 霍建华赵丽颖演绎绝美爱情