首页 > 代码库 > 数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
上一篇提到,那份源码的使用是有限制的。
这次来完善一下。其实就是迭代多次,使得最后一次刚好在一个线程块可以求和。
完善部分:
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=262144 const int num_elements = 1<<20; // generate random input on the host std::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千万),所以先这样。
数组求和的快速方法(利用cuda的共享内存)--第二部分之程序完善
声明:以上内容来自用户投稿及互联网公开渠道收集整理发布,本网站不拥有所有权,未作人工编辑处理,也不承担相关法律责任,若内容有误或涉及侵权可进行投诉: 投诉/举报 工作人员会在5个工作日内联系你,一经查实,本站将立刻删除涉嫌侵权内容。