存在错误的程序是sample中的scanLargeArray,在指出具体错误之前,让我们先看一下这个程序的运行结果吧.
我的本本是130M,运行结果如下:
Running parallel prefix sum (prescan) of 1000000 elements
This version is work efficient (O(n) adds)
and has very few shared memory bank conflicts
Test PASSED
Average GPU execution time: 0.082135 ms
CPU execution time: 8.225319 ms
Check out the CUDA Data Parallel Primitives Library for more on scan.
http://www.gpgpu.org/developer/cudpp
Press ENTER to exit...
哇,GPU比CPU快100倍!GPU真有前途....拍手!
等会,记得好像prefix sum也是一个受带宽限制的任务吧,我的本本上的GPU有那么强壮吗?
怀着好奇心,计算了一下,就算至少扫描一遍数据:
一个float是4字节,4*1000000/0.082135*1000=48.7GB/s !!!
记得130M只有128bit的带宽啊,有那么厉害吗?
用GPU-Z一查,晕了,才25.6GB/s。(见下图)
肯定有问题!还好,有source可查,看了好几遍source,还是一下看不出问题。
最后,在测量时间的地方发现了问题所在。
请看scanLargeArray.cu中的下面一段程序:
preallocBlockSums(num_elements);
// run once to remove startup overhead
prescanArray(d_odata, d_idata, num_elements);
// Run the prescan
cutStartTimer(timerGPU);
for (unsigned int i = 0; i < num_test_iterations; i++)
{
//printf("prescanArray/n");
prescanArray(d_odata, d_idata, num_elements);
}
cutStopTimer(timerGPU);
deallocBlockSums();
// copy result from device to host
cutilSafeCall(cudaMemcpy( h_data, d_odata, sizeof(float) * num_elements,
cudaMemcpyDeviceToHost));
看出问题吗?哈哈,其实bug很低级,就是测试的都是内核的启动时间而非执行时间!
整个source中没有任何cudaThreadSynchronize!最后的同步是在
cutilSafeCall(cudaMemcpy( h_data, d_odata, sizeof(float) * num_elements,
cudaMemcpyDeviceToHost));
这里,而这时,计时已经结束了。
为证实这点,小小改了一下程序,在计时结束前,加上一句cudaThreadSynchronize调用。如下:
// Run the prescan
cutStartTimer(timerGPU);
for (unsigned int i = 0; i < num_test_iterations; i++)
{
//printf("prescanArray/n");
prescanArray(d_odata, d_idata, num_elements);
}
cudaThreadSynchronize();// add by l7331014
cutStopTimer(timerGPU);
这下运行结果变成了如下数据:
Running parallel prefix sum (prescan) of 1000000 elements
This version is work efficient (O(n) adds)
and has very few shared memory bank conflicts
Test PASSED
Average GPU execution time: 3.330053 ms
CPU execution time: 9.927999 ms
Check out the CUDA Data Parallel Primitives Library for more on scan.
http://www.gpgpu.org/developer/cudpp
Press ENTER to exit...
再计算一次带宽:
4*1000000/3.330053*1000=1.2GB/s !
这个才是真正的性能数据。真是天上地下啊,呵呵。
又运行了自己写的一个测试内核启动时间的小程序,结果如下:
Start time test .....
Processing time: 0.088433 (ms)
Finished !
这下算证实了。
最后说一下,sdk sample中的scan算法不是好的算法,在这里就不多提了,我会在以后的文章中分析为何不好。呵呵。