-个例子

上一节，已经简单的说了一下CUDA C的基本语法；因而在本节，兄弟决定以一个例子为基础说明CUDA程序的基本组成部分，不过说实话兄弟选择的例子并不太好，这个例子就是采用积分法计算圆周率/π的值。其计算原理是：在[0,1]范围内积分1/(1+x*x)

首先，让我们看一下在CPU上的计算流程，其计算流程如下

/*

串行计算PI的程序，基本思想为：将积分区间均分为num小块，将每小块的面积加起来。

*/ 

   float cpuPI(int num){

      float sum=0.0f;

      float temp;

      for(int i=0;i<num;i++){

          temp=(i+0.5f)/num;

      // printf("%f\n",temp);

          sum+=4/(1+temp*temp);

      // printf("%f\n",sum);

      }

      return sum/num;

   }

很明显，我们可以将for循环分解，使用CUDA处理。

有一个问题就是：for内部对sum变量的更新是互斥的，而CUDA中并没有浮点原子函数，对于这个问题的解决方案是：将for循环内部的两个语句拆开，分成两个内核函数来做运算。内核函数英文名为kernel，就是一个能够在GPU上运算的模块。第一个内核计算计算每个小积分块面积，并将每个block内所有线程对应的积分块面积加起来，存入全局存储器；第二个内核将前一个内核存入的数据加起来。下面是kernel代码：

/*

在GPU上计算PI的程序，要求块数和块内线程数都是2的幂

前一部分为计算block内归约，最后大小为块数

后一部分为单个block归约，最后存储到*pi中。

*/

/*

在GPU上计算PI的程序，要求块数和块内线程数都是2的幂

前一部分为计算block内归约，最后大小为块数

后一部分为单个block归约，最后存储到*pi中。

*/

   __global__ void reducePI1(float *d_sum,int num){

   int id=blockIdx.x*blockDim.x+threadIdx.x;//线程索引

   int gid=id;

   float temp;

   extern float __shared__ s_pi[];//动态分配，长度为block线程数

   s_pi[threadIdx.x]=0.0f;

   while(gid<num){

      temp=(gid+0.5f)/num;//当前x值

      s_pi[threadIdx.x]+=4.0f/(1+temp*temp);

      gid+=blockDim.x*gridDim.x;

   }

   for(int i=(blockDim.x>>1);i>0;i>>=1){

      if(threadIdx.x<i){

          s_pi[threadIdx.x]+=s_pi[threadIdx.x+i];

      }

      __syncthreads();

   }

   if(threadIdx.x==0)

   d_sum[blockIdx.x]=s_pi[0];

   }

   __global__ void reducePI2(float *d_sum,int num,float *d_pi){

   int id=threadIdx.x;

   extern float __shared__ s_sum[];

   s_sum[id]=d_sum[id];

   __syncthreads();

   for(int i=(blockDim.x>>1);i>0;i>>=1){

      if(id<i)

      s_sum[id]+=s_sum[id+i];

      __syncthreads();

   }

// printf("%d,%f\n",id,s_sum[id]);

   if(id==0){

   *d_pi=s_sum[0]/num;

// printf("%d,%f\n",id,*pi);

   }

   }

其中__syncthreads()是CUDA的内置命令，其作用是保证block内的所有线程都已经运行到调用__syncthreads()的位置。

   由上面的代码可以看出，使用使用CUDA的主要阻碍在于数据相关性。

   一般而言，CUDA程序的基本模式是：

一、分配内存空间和显存空间

二、初始化内存空间

三、将要计算的数据从内存上复制到显存上

四、执行kernel计算

五、将计算后显存上的数据复制到内存上

六、处理复制到内存上的数据

这个程序在我的机器(CPU 2.0GHZ,GPU GTX295)上的加速比超过100，不知道在你们的机器上能够加速多少？