-个例子-CSDN博客

Normal 0 7.8 磅 0 2 false false false EN-US ZH-CN X-NONE

上一节，已经简单的说了一下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

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;