數組求和的快速方法（利用cuda的共享內存）--第二部分之程序完善

上一篇提到，那份源碼的使用是有限制的。

這次來完善一下。其實就是迭代多次，使得最後一次剛好在一個線程塊可以求和。

完善部分：

template
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<<>>(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 h_input(num_elements);
	for(int i = 0; i < h_input.size(); ++i)
	{
		h_input[i] = 1;//random_num();
	}

	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_blocks.

size_t num_blocks = (num_elements/block_size) + ((num_elements%block_size) ? 1 : 0);

萬一第一次求出的num_blocks大於線程塊的最大數量，一般是65535，那就不行了。

所以如果數組的元素數量大於1024*65535，那就無法計算了。

解決這中問題的通常方法，是讓一個線程串行執行多個相同任務。

由於求解的問題暫時沒有超過這個數量級（6-7千萬），所以先這樣。