程序员人生网站导航

【CUDA并行编程之七】数组元素之和

栏目：服务器时间：2015-01-31 10:19:05

现在需要求得1个数组的所有元素之和，之前感觉似乎不太可能，由于每一个线程只处理1个元素，没法将所有元素联系起来，但是最近学习了1段代码可以实现，同时也对shared memory有了进1步的理解。

1、C++串行实现

串行实现的方法非常之简单，只要将所有元素顺次相加就可以够得到相应的结果，实际上我们重视的不是结果，而是运行的效力。那末代码以下：

array_sum.cc:

#include<iostream> #include<stdio.h> #include "kmeans.h" using namespace std; const int cnt = 100000; int main() { int *a = new int[cnt]; for(int i=0;i<cnt;i++) { a[i] = i+1; } double t = wtime(); for(int i=0;i<cnt;i++) sum += a[i]; printf("computation elapsed %.8f ",wtime()-t); return 0; }

wtime.cu:

#include <sys/time.h> #include <stdio.h> #include <stdlib.h> double wtime(void) { double now_time; struct timeval etstart; struct timezone tzp; if (gettimeofday(&etstart, &tzp) == ⑴) perror("Error: calling gettimeofday() not successful. "); now_time = ((double)etstart.tv_sec) + /* in seconds */ ((double)etstart.tv_usec) / 1000000.0; /* in microseconds */ return now_time; }

运行结果：

2、CUDA并行实现

先上代码然后再进行解释：

#include <iostream> #include <stdio.h> #include "kmeans.h" using namespace std; const int count = 1000; void generate_data(int *arr) { for(int i=0;i<count;i++) { arr[i] = i+1; } } int nextPowerOfTwo(int n) { n--; n = n >> 1 | n; n = n >> 2 | n; n = n >> 4 | n; n = n >> 8 | n; n = n >> 16 | n; //n = n >> 32 | n; //For 64-bits int return ++n; } /* cnt : count cnt2 : next power of two of count */ __global__ static void compute_sum(int *array,int cnt , int cnt2) { extern __shared__ unsigned int sharedMem[]; sharedMem[threadIdx.x] = (threadIdx.x < cnt) ? array[threadIdx.x] : 0 ; __syncthreads(); //cnt2 "must" be a power of two! for( unsigned int s = cnt2/2 ; s > 0 ; s>>=1 ) { if( threadIdx.x < s ) { sharedMem[threadIdx.x] += sharedMem[threadIdx.x + s]; } __syncthreads(); } if(threadIdx.x == 0) { array[0] = sharedMem[0]; } } int main() { int *a = new int[count]; generate_data(a); int *deviceArray; cudaMalloc( &deviceArray,count*sizeof(int) ); cudaMemcpy( deviceArray,a,count*sizeof(int),cudaMemcpyHostToDevice ); int npt_count = nextPowerOfTwo(count);//next power of two of count //cout<<"npt_count = "<<npt_count<<endl; int blockSharedDataSize = npt_count * sizeof(int); double t = wtime(); for(int i=0;i<count;i++) { compute_sum<<<1,count,blockSharedDataSize>>>(deviceArray,count,npt_count); } printf("computation elapsed %.8f ",wtime()-t); int sum ; cudaMemcpy( &sum,deviceArray,sizeof(int),cudaMemcpyDeviceToHost ); cout<<"sum = "<<sum<<endl; return 0; }

主函数：
line58：为数组a赋初值，维度为count。

line60~62：定义device变量并分配内存，将数组a的值拷贝到显存上去。

line63：nextPowerOfTwo是非常精巧的1段代码，它计算大于等于输入参数n的第1个2的幂次数。至于为何这么做要到kernel函数里面才能明白。

line68：compute_sum中的"1"为block的数量，"count"为每一个block里面的线程数，"blockSharedDataSize"为同享内存的大小。

核函数compute_sum：

line35：定义shared memory变量。

line36：将threadIdx.x小于cnt的对应的sharedMem的内存区赋值为数组array中的值。

line39~47：这段代码的功能就是将所有值求和以后放到了shareMem[0]这个位置上。这段代码要好好体会1下，它把本来计算复杂度为O(n)的串行实现的时间效力通过并行到达了O(logn)。最后将结果保存到array[0]并拷贝回主存。

makefile：

cu: nvcc cuda_array_sum.cu wtime.cu ./a.out

结果：