CUDA实现平衡树求前缀和

应实验要求，使用CUDA实现平衡树的前缀求和算法，很可惜只是做了一个简单的版本，没有时间做更多的优化，在此记录下来以后在进行深入优化。

平衡树求前缀和的思想是使用逻辑上的树形结构，先一次遍历所有叶子节点实现相邻的两个节点相加，层层上播直到根节点。然后从根节点在执行一次向下的播送过程，最后所有的前缀和都是存放在叶子节点中。

为了简单起见使用一维数组保存所有的叶子节点和相应的中间节点，然后使用一些小标小技术让其满足逻辑上的树形结构。

插入串行代码：

int GetJump(int n,int step){int i;int sum = 0;if(step == 0)return 0;for(i = 1;i <= step;i *= 2)sum += n / i;return sum;}void PrefixST(int *input,int n){int h,k,a,b;int broadcastA,broadcastB;int step = 1;/*   Up Sum*/for(h = n / 2;h > 0;h /= 2){for(k = 0;k < h;k++){a = 2 * k + GetJump(n,step / 2);b = 2 * k + 1 + GetJump(n,step / 2);input[GetJump(n,step) + k] = input[a] + input[b];}step *= 2;}/*   Down Broadcast*///printf("step : %d",step);step /= 2;for(h = 1;h < n;h *= 2){for(k = 0;k < h;k++){if(k == 0){broadcastA = input[k + GetJump(n,step)];a = 2 * k + 1 + GetJump(n,step / 2);input[a] = broadcastA;}else {broadcastA = input[k + GetJump(n,step)];broadcastB = input[k + GetJump(n,step) - 1];a = 2 * k + GetJump(n,step / 2);b = 2 * k + 1 + GetJump(n,step / 2);input[a] += broadcastB;input[b] = broadcastA;}}step /= 2;}}

然后依据这个串行代码做一些小小的改进就能得到CUDA的代码，其中把循环遍历改为线程的tid就行。这里使用最简单的1 blocks，n 列threads的方式，为了与一维数组进行很好的关联，便于操作线程。

插入并行代码：

__device__ int GetJump(int n,int step){int i,sum = 0;if(step == 0)return 0;for(i = 1;i <= step;i *= 2) //n + n / 2 + n / 4 +......sum += n / i;return sum;}__global__ void PrefixST(int *input,int *output,int n){__shared__ int temp[Num * 2]; int h;int step;int a,b;int broadcastA,broadcastB; int tid = threadIdx.x;temp[tid] = input[tid];temp[tid + n] = 0;/* Up Sum */step = 1;for(h = n / 2;h > 0;h = h / 2){__syncthreads();if(tid < h){a = (step != 1) * GetJump(n,step / 2) + 2 * tid;b = (step != 1) * GetJump(n,step / 2) + 2 * tid + 1;temp[tid + GetJump(n,step)] = temp[a] + temp[b];}step *= 2;}/* Down Broadcast */__syncthreads();step = step / 2;for(h = 1;h < n;h *= 2){if(tid == 0){// Most left childbroadcastA = temp[tid + GetJump(n,step)];//fathera = 2 * tid + 1 + GetJump(n,step / 2);//right childtemp[a] = broadcastA;}else if(tid < h && tid > 0){broadcastA = temp[tid + GetJump(n,step)];//fatherbroadcastB = temp[tid + GetJump(n,step) - 1];//father's left brothera = 2 * tid + GetJump(n,step / 2);//left childb = 2 * tid + 1 + GetJump(n,step / 2);//right child temp[a] += broadcastB; temp[b] = broadcastA;}step /= 2;}__syncthreads();output[tid] = temp[tid];output[tid + n] = temp[tid + n];} 

使用CUDA的自带计时器：发现时间效果很差，并行的算出来还没得串行的好，看来还有很大的优化空间...

float time;cudaEvent_t start,stop;cudaEventCreate(&start);cudaEventCreate(&stop);cudaEventRecord(start,0);/*Function runs*/cudaEventRecord(stop,0);cudaEventSynchronize(stop);cudaEventElapsedTime(&time,start,stop);