CUDA程序优化小记（五）

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要循环2¹³次进行加的操作。更糟糕的是，这一部分工作和GPU上的求和还是串行的，无法真实反映GPU完成的任务。于是这回将主要精力放在如何将这一部分代码放在GPU上运行。

         CUDA有一大特色，那就是共享存储器。共享存储器位于每个多处理器内，是片上存储器，它的作用域在一个块（BLOCK）内。于是在块内进行各个线程的结果求和，比较符合我们的期望，也就是说，线程执行y=X²这样的操作，将结果放在共享存储器中，再选择一个线程执行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( );// 选择任意一个线程进行计时
 
    //