之前学习过omp编程,结合现在的cuda,我觉得要理解并行编程,首先理解 划分和 规约这两个概念。也许你的算法学的更加扎实。划分是《算法》里面的一个重要思想,将一个大的问题或任务, 分解成小问题小任务,各个击破,最后 归并结果;规约是《cuda**》书上介绍的一个入门的重要思想,规约算法(reduction)用来求连加、连乘、最值等,应用广泛。每次循环参加运算的线程减少一半。
不管算法的思想如何花样,万变不离其中的一点--将一个大的任务分解成小的任务集合,分解原则是粒度合适尽量小、数据相关性尽量小。如此而已。因为,我们用GPU是为了加速,要加速必须提高执行任务的并行度!明白这个道理,那么我们将绞尽脑汁地去想方设法分析自己手上的任务,分解、分解、分解!
这里拿规约来说事情,因为,规约这个东西,似乎可以拿来单做9*9乘法表来熟悉,熟悉了基础的口诀,那么99*99的难题也会迎刃而解。
ex:矢量加法,要实现N=64*256长度的矢量的累加和。假设a+b计算一次耗时t。
cpu计算:
显然单核的话需要64*256*t。我们容忍不了。
gpu计算:
最初的设想,我们如果有个gpu能同时跑N/2个线程,我们这N/2个线程同时跑,那么不是只需要t时间就能将N个数相加编程N/2个数相加了吗?对的。这一轮我们用了t时间;
接着的想法,我们不断的递归这个过程,能发现吗?第二轮,我们用N/2/2个线程同时跑,剩下N/2/2个数相加,这一轮我们同样用了t时间;
一直这样想下去吧,最后一轮,我们用了1个线程跑,剩下1个数啦,这就是我们的结果!!!
每一轮时间都为t,那么理想情况,我们用了多少轮这样的计算呢?
计算次数=log(N)=6*8=48,对的,只用了48轮,也就是说,我们花了48*t的时间!
规约就是这样,很简单,很好用,我们且不管程序后期的优化,单从这个算法分析上来说,从时间复杂度N降到了logN,这在常规算法上,要提高成这样的效率,是不得了的,
这是指数级别的效率提高!所以,你会发现,GPU有CPU无法取代的得天独厚的优势————处理单元真心多啊!
规约求和的核函数代码如下:
__global__ void RowSum(float* A, float* B){
int bid = blockIdx.x;
int tid = threadIdx.x;
__shared__ s_data[128];
//read data to shared memory
s_data[tid] = A[bid*128 + tid];
__synctheads(); //sync
for(int i=64; i>0; i/=2){
if(tid<i)
s_data[tid] = s_data[tid] + s_data[tid+i] ;
__synctheads();
}
if(tid==0)
B[bid] = s_data[0];
}