CUDA程序优化小记(五)
CUDA全称Computer Unified Device Architecture(计算机同一设备架构),它的引入为计算机计算速度质的提升提供了可能,从此微型计算机也能有与大型机相当计算的能力。可是不恰当地使用CUDA技术,不仅不会让应用程序获得提升,反而会比普通CPU的计算还要慢。最近我通过学习《GPGPU编程技术》这本书,深刻地体会到了这一点,并且用CUDARuntime应用改写书上的例子程序;来体会CUDA技术给我们计算能力带来的提升。
原创文章,反对未声明的引用。原博客地址:http://blog.csdn.net/gamesdev/article/details/18140163
上一版程序以非常惊人的速度完成了数据求和的计算,不过在GPU计算各个分组之后,CPU还要将各个处理器的和再求一遍和。也就是说,在CPU上运行了这样一段代码:
// 在CPU中将结果组合起来
int totalResult;
for ( int i = 0; i <THREAD_NUM * BLOCK_NUM; ++i )
{
totalResult += pResult[i];
}
这样等于是将GPU的负担给了一部分给CPU做,在这里,定义BLOCK_NUM和THREAD_NUM分别为32和256,也就是说CPU要循环213次进行加的操作。更糟糕的是,这一部分工作和GPU上的求和还是串行的,无法真实反映GPU完成的任务。于是这回将主要精力放在如何将这一部分代码放在GPU上运行。
CUDA有一大特色,那就是共享存储器。共享存储器位于每个多处理器内,是片上存储器,它的作用域在一个块(BLOCK)内。于是在块内进行各个线程的结果求和,比较符合我们的期望,也就是说,线程执行y=X2这样的操作,将结果放在共享存储器中,再选择一个线程执行z=∑y的操作,最后将w=∑z交给CPU来执行。下面是修改后的内核代码:
__global__ static voidKernel_SquareSum( int* pIn, size_t* pDataSize,
int*pOut, clock_t* pTime )
{
// 声明一个动态分配的共享存储器
extern __shared__ int sharedData[];
const size_t computeSize =*pDataSize / THREAD_NUM;
const size_t tID = size_t(threadIdx.x );// 线程
const size_t bID = size_t(blockIdx.x );// 块
// 开始计时
if ( tID == 0 ) pTime[bID] =clock( );// 选择任意一个线程进行计时
//