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

柔光的暖阳◎ 2022-08-01 12:17 179阅读 0赞

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

**一、C++串行实现**

串行实现的方法非常之简单，只要将所有元素依次相加就能够得到相应的结果，实际上我们注重的不是结果，而是运行的效率。那么代码如下：

**array\_sum.cc:**

**\[cpp\]**  [view plain][] [copy][view plain] [![在CODE上查看代码片][CODE]][CODE_CODE] [![派生到我的代码片][ico_fork.svg]][ico_fork.svg 1]

1.  \#include<iostream>  
2.  \#include<stdio.h>  
3.  \#include "kmeans.h"  
4.    
5.  **using** **namespace** std;  
6.    
7.  **const** **int** cnt = 100000;  
8.    
9.  **int** main()  
10. \{  
11.     **int** \*a = **new** **int**\[cnt\];  
12.   
13.     **for**(**int** i=0;i<cnt;i++)  
14.     \{  
15.         a\[i\] = i+1;   
16.     \}  
17.   
18.     **double** t = wtime();  
19.       
20.     **for**(**int** i=0;i<cnt;i++)  
21.         sum += a\[i\];  
22.       
23.     printf("computation elapsed %.8f \\n",wtime()-t);  
24.   
25.     **return** 0;  
26. \}

**wtime.cu:**

**\[cpp\]**  [view plain][] [copy][view plain] [![在CODE上查看代码片][CODE]][CODE_CODE] [![派生到我的代码片][ico_fork.svg]][ico_fork.svg 1]

1.  \#include <sys/time.h>  
2.  \#include <stdio.h>  
3.  \#include <stdlib.h>  
4.    
5.  **double** wtime(**void**)   
6.  \{  
7.      **double**          now\_time;  
8.      **struct** timeval  etstart;  
9.      **struct** timezone tzp;  
10.   
11.     **if** (gettimeofday(&etstart, &tzp) == -1)  
12.         perror("Error: calling gettimeofday() not successful.\\n");  
13.   
14.     now\_time = ((**double**)etstart.tv\_sec) +              /\* in seconds \*/  
15.                ((**double**)etstart.tv\_usec) / 1000000.0;  /\* in microseconds \*/  
16.     **return** now\_time;  
17. \}

运行结果：

![SouthEast][]

**二、CUDA并行实现**

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

**\[cpp\]**  [view plain][] [copy][view plain] [![在CODE上查看代码片][CODE]][CODE_CODE] [![派生到我的代码片][ico_fork.svg]][ico_fork.svg 1]

1.  \#include <iostream>  
2.  \#include <stdio.h>  
3.  \#include "kmeans.h"  
4.    
5.  **using** **namespace** std;  
6.    
7.  **const** **int** count = 1000;  
8.    
9.  **void** generate\_data(**int** \*arr)  
10. \{  
11.     **for**(**int** i=0;i<count;i++)   
12.     \{  
13.         arr\[i\] = i+1;  
14.     \}  
15. \}  
16.   
17. **int** nextPowerOfTwo(**int** n)  
18. \{  
19.     n--;  
20.     n = n >> 1 | n;  
21.     n = n >> 2 | n;  
22.     n = n >> 4 | n;  
23.     n = n >> 8 | n;  
24.     n = n >> 16 | n;  
25.     //n = n >> 32 | n; //For 64-bits int   
26.     **return** ++n;  
27. \}  
28.   
29. /\* 
30. cnt : count  
31. cnt2 : next power of two of count  
32. \*/  
33. \_\_global\_\_ **static** **void** compute\_sum(**int** \*array,**int** cnt , **int** cnt2)  
34. \{  
35.     **extern** \_\_shared\_\_ unsigned **int** sharedMem\[\];  
36.     sharedMem\[threadIdx.x\] = (threadIdx.x < cnt) ? array\[threadIdx.x\] : 0 ;  
37.     \_\_syncthreads();  
38.   
39.     //cnt2 "must" be a power of two!  
40.     **for**( unsigned **int** s = cnt2/2 ; s > 0 ; s>>=1 )  
41.     \{  
42.         **if**( threadIdx.x < s )      
43.         \{  
44.             sharedMem\[threadIdx.x\] += sharedMem\[threadIdx.x + s\];  
45.         \}  
46.         \_\_syncthreads();  
47.     \}  
48.     **if**(threadIdx.x == 0)  
49.     \{  
50.         array\[0\] = sharedMem\[0\];      
51.     \}  
52. \}  
53.   
54.   
55. **int** main()  
56. \{  
57.     **int** \*a = **new** **int**\[count\];  
58.     generate\_data(a);  
59.   
60.     **int** \*deviceArray;  
61.     cudaMalloc( &deviceArray,count\***sizeof**(**int**) );  
62.     cudaMemcpy( deviceArray,a,count\***sizeof**(**int**),cudaMemcpyHostToDevice );  
63.     **int** npt\_count = nextPowerOfTwo(count);//next power of two of count  
64.     //cout<<"npt\_count = "<<npt\_count<<endl;  
65.     **int** blockSharedDataSize = npt\_count \* **sizeof**(**int**);  
66.   
67.     **double** t = wtime();  
68.     **for**(**int** i=0;i<count;i++)  
69.     \{  
70.         compute\_sum<<<1,count,blockSharedDataSize>>>(deviceArray,count,npt\_count);  
71.     \}  
72.     printf("computation elapsed %.8f \\n",wtime()-t);  
73.   
74.     **int** sum ;  
75.     cudaMemcpy( &sum,deviceArray,**sizeof**(**int**),cudaMemcpyDeviceToHost );  
76.     cout<<"sum = "<<sum<<endl;  
77.       
78.     **return** 0;  
79. \}

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

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

**line63：**nextPowerOfTwo是非常精妙的一段代码，它计算大于等于输入参数n的第一个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\]这个位置上。这段代码要好好体会一下，它把**原本计算复杂度为O(n)的串行实现的时间效率通过并行达到了O(logn)**。最后将结果保存到array\[0\]并拷贝回主存。

makefile：

**\[cpp\]**  [view plain][] [copy][view plain] [![在CODE上查看代码片][CODE]][CODE_CODE] [![派生到我的代码片][ico_fork.svg]][ico_fork.svg 1]

1.  cu:  
2.      nvcc cuda\_array\_sum.cu wtime.cu  
3.      ./a.out

结果：

![SouthEast 1][]

**三、效率对比**

我们通过修改count的值并且加大循环次数来观察变量的效率的差别。

代码修改：

![SouthEast 2][]

![SouthEast 3][]

![SouthEast 4][]

**运行时间对比：**

<table style="border-collapse:collapse; margin-top:10px; margin-bottom:10px; width:600px; background-color:inherit"> 
 <tbody style="background-color:inherit"> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <strong><span style="font-size:18px">count</span></strong> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> <strong><span style="font-size:18px">串行（s）</span></strong></td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <strong><span style="font-size:18px">并行（s）</span></strong> 
    </div> </td> 
  </tr> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:27px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">1000</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:27px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">0.00627995</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:27px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">0.00345612</span> 
    </div> </td> 
  </tr> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">10000</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">0.29315591</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">0.06507015</span> 
    </div> </td> 
  </tr> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">100000</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">25.18921304</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">0.65188980</span> 
    </div> </td> 
  </tr> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:26px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">1000000</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:26px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">2507.66827798</span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:26px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px">5.61648989</span> 
    </div> </td> 
  </tr> 
  <tr style="background-color:inherit"> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px"><br style="background-color:inherit"> </span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px"><br style="background-color:inherit"> </span> 
    </div> </td> 
   <td style="word-break:break-all; border:1px solid rgb(153,153,153); padding:5px 10px; min-height:25px; min-width:25px; height:25px; background-color:inherit; width:179px"> 
    <div style="min-width:2px; background-color:inherit"> 
     <span style="font-size:18px"><br style="background-color:inherit"> </span> 
    </div> </td> 
  </tr> 
 </tbody> 
</table>

哈哈，可见在数据量大的情况下效率还是相当不错的。

Author：忆之独秀

Email：leaguenew@qq.com

注明出处：[http://blog.csdn.net/lavorange/article/details/43031419][http_blog.csdn.net_lavorange_article_details_43031419]

[view plain]: http://blog.csdn.net/lavorange/article/details/43031419#
[CODE]: https://code.csdn.net/assets/CODE_ico.png
[CODE_CODE]: https://code.csdn.net/snippets/586751
[ico_fork.svg]: https://code.csdn.net/assets/ico_fork.svg
[ico_fork.svg 1]: https://code.csdn.net/snippets/586751/fork
[SouthEast]: /images/20220731/624bfe5de68e4381b3850d4c34925f59.png
[SouthEast 1]: /images/20220731/d84d758fb6d24f2dabfedf2b7e87ae1d.png
[SouthEast 2]: /images/20220731/f051a7497beb437abef2f91551421081.png
[SouthEast 3]: /images/20220731/993714bb54064d9da46afae507a1376f.png
[SouthEast 4]: /images/20220731/0257ea78ba39484fb9d53ebeb533119a.png
[http_blog.csdn.net_lavorange_article_details_43031419]: http://blog.csdn.net/lavorange/article/details/43031419