数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
来源:互联网 发布:win7固态硬盘优化 编辑:程序博客网 时间:2024/06/05 16:47
上一篇提到,那份源码的使用是有限制的。
这次来完善一下。其实就是迭代多次,使得最后一次刚好在一个线程块可以求和。
完善部分:
template<class DType>DType array_sum_gpu(DType *dev_array,const int array_size,DType *dev_result){//const size_t max_block_size = 512;//目前有些gpu的线程块最大为512,有些为1024.const size_t block_size = 512;//线程块的大小。size_t num_elements = array_size;size_t num_blocks = (num_elements/block_size) + ((num_elements%block_size) ? 1 : 0);double *dev_input_array = 0;double *dev_block_sums = 0;//一个线程块一个和。while(num_elements > block_size){if(dev_block_sums == 0)//第一次{dev_input_array = dev_array;}else //除了第一次{if(dev_input_array != dev_array)cudaFree(dev_input_array);dev_input_array = dev_block_sums;}num_blocks = (num_elements/block_size) + ((num_elements%block_size) ? 1 : 0);//给输出结果分配内存cudaMalloc((void**)&dev_block_sums, sizeof(double) * (num_blocks ));// launch one kernel to compute, per-block, a partial sum//把每个线程块的和求出来block_sum<<<num_blocks,block_size,block_size * sizeof(double)>>>(dev_input_array, dev_block_sums, num_elements);num_elements = num_blocks;}block_sum<<<1,num_elements,num_elements * sizeof(double)>>>(dev_block_sums, dev_result, num_elements);double result = 0;cudaMemcpy(&result, dev_result, sizeof(double), cudaMemcpyDeviceToHost);cudaFree(dev_block_sums);return result;}
核函数block_sum还是原来的代码。
下面是测试我的代码;
void test_sum2(){// create array of 256k elements//const int num_elements = 1<<18;//=512*512=262144const int num_elements = 1<<20;// generate random input on the hoststd::vector<double> h_input(num_elements);for(int i = 0; i < h_input.size(); ++i){h_input[i] = 1;//random_num<double>();}const double host_result = std::accumulate(h_input.begin(), h_input.end(), 0.0f);std::cerr << "Host sum: " << host_result << std::endl;// move input to device memory//分配内存double *d_input = 0;cudaMalloc((void**)&d_input, sizeof(double) * num_elements);cudaMemcpy(d_input, &h_input[0], sizeof(double) * num_elements, cudaMemcpyHostToDevice);double *dev_result=0;cudaMalloc((void**)&dev_result,sizeof(double));double sum = array_sum_gpu(d_input,num_elements,dev_result);std::cout << "Device sum: " << sum << std::endl;}
可以把数组数量num_elements调到很大,代码仍然能运行正确。但是如果是之前那份代码,就不可以了。
其实这个程序还是有点限制的。
请注意第一次求num_blocks.
size_t num_blocks = (num_elements/block_size) + ((num_elements%block_size) ? 1 : 0);
万一第一次求出的num_blocks大于线程块的最大数量,一般是65535,那就不行了。
所以如果数组的元素数量大于1024*65535,那就无法计算了。
解决这中问题的通常方法,是让一个线程串行执行多个相同任务。
由于求解的问题暂时没有超过这个数量级(6-7千万),所以先这样。
1 0
- 数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
- 数组求和的快速方法(利用cuda的共享内存)--第一部分之源码分析
- 数组求和的快速方法(利用cuda的共享内存)--第三部分之性能分析
- 快速开发CUDA程序的方法
- CUDA 共享内存的动态分配
- 数组求和的方法
- javaSE_第二部分 数组及内存浅析_方法的重载
- CUDA编程(七)共享内存与Thread的同步
- CUDA编程(七)共享内存与Thread的同步
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- (译)如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- 如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- 如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- 如何优化cocos2d程序的内存使用和程序大小:第二部分(完)
- 网曝小伙手机贴膜月入3万 引发网友质疑 新辰深入调查真相
- JTextpane 添加行号
- Hadoop架构设计、运行原理详解
- linux重定向
- 练习题(1)
- 数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
- Handler消息传递机制
- 使用CocoaPods管理依赖库
- 虚拟机有分区的情况下,扩磁盘
- 如何重写hashcode方法
- sqlite之我见--C/C++ API接口介绍 .
- 改变世界的九大算法
- Linux下文本搜索工具grep命令使用入门
- Websocket java api (jsr 356) study notes