为什么要使用共享内存呢，因为共享内存的访问速度快。这是首先要明确的，下面详细研究。
cuda程序中的内存使用分为主机内存（host memory） 和 设备内存（device memory），我们在这里关注的是设备内存。设备内存都位于gpu之上，前面我们看到在计算开始之前，每次我们都要在device上申请内存空间，然后把host上的数据传入device内存。cudaMalloc（）申请的内存，还有在核函数中用正常方法申请的变量的内存。这些内存叫做全局内存，那么还有没有别的内存种类呢？常用的还有共享内存，常量内存，纹理内存，他们都用一些不正常的方法申请。
他们的申请方法如下：
共享内存：__shared__ 变量类型 变量名；
常量内存：__constant__ 变量类型 变量名；
纹理内存：texture<变量类型> 变量名；
| 存储类型 | 寄存器 | 共享内存 | 纹理内存 | 常量内存 | 全局内存 |
| 带宽 | ～8TB/s | ～1.5TB/s | ～200MB/s | ～200MB/s |
～200MB/s |
| 延迟 | 1个周期 | 1～32周期 | 400～600周期 | 400～600周期 | 400～600周期 |
他们在不同的情况下有各自的作用，他们最大的区别就是带宽不同，通俗说就是访问速度不同。后面三个看起来没什么不同，但是他们在物理结构方面有差别，适用于不同的情况。
共享内存实际上是可受用户控制的一级缓存。申请共享内存后，其内容在每一个用到的block被复制一遍，使得在每个block内，每一个thread都可以访问和操作这块内存，而无法访问其他block内的共享内存。这种机制就使得一个block之内的所有线程可以互相交流和合作。下面的例子中就显示了线程之间的交流和合作。
这个例子计算的是两个向量的点积。
1 /*
2 * Copyright 1993-2010 NVIDIA Corporation. All rights reserved.
3 *
4 * NVIDIA Corporation and its licensors retain all intellectual property and
5 * proprietary rights in and to this software and related documentation.
6 * Any use, reproduction, disclosure, or distribution of this software
7 * and related documentation without an express license agreement from
8 * NVIDIA Corporation is strictly prohibited.
9 *
10 * Please refer to the applicable NVIDIA end user license agreement (EULA)
11 * associated with this source code for terms and conditions that govern
12 * your use of this NVIDIA software.
13 *
14 */
15
16
17 #include "../common/book.h"
18
19 #define imin(a,b) (a<b?a:b)
20
21 const int N = 33 * 1024;
22 const int threadsPerBlock = 256;
23 const int blocksPerGrid =
24 imin( 32, (N+threadsPerBlock-1) / threadsPerBlock );
25
26
27 __global__ void dot( float *a, float *b, float *c ) {
28 __shared__ float cache[threadsPerBlock];
29 int tid = threadIdx.x + blockIdx.x * blockDim.x;
30 int cacheIndex = threadIdx.x;
31
32 float temp = 0;
33 while (tid < N) {
34 temp += a[tid] * b[tid];
35 tid += blockDim.x * gridDim.x;
36 }
37
38 // set the cache values
39 cache[cacheIndex] = temp;
40
41 // synchronize threads in this block
42 __syncthreads();
43
44 // for reductions, threadsPerBlock must be a power of 2
45 // because of the following code
46 int i = blockDim.x/2;
47 while (i != 0) {
48 if (cacheIndex < i)
49 cache[cacheIndex] += cache[cacheIndex + i];
50 __syncthreads();
51 i /= 2;
52 }
53
54 if (cacheIndex == 0)
55 c[blockIdx.x] = cache[0];
56 }
57
58
59 int main( void ) {
60 float *a, *b, c, *partial_c;
61 float *dev_a, *dev_b, *dev_partial_c;
62
63 // allocate memory on the cpu side
64 a = (float*)malloc( N*sizeof(float) );
65 b = (float*)malloc( N*sizeof(float) );
66 partial_c = (float*)malloc( blocksPerGrid*sizeof(float) );
67
68 // allocate the memory on the GPU
69 HANDLE_ERROR( cudaMalloc( (void**)&dev_a,
70 N*sizeof(float) ) );
71 HANDLE_ERROR( cudaMalloc( (void**)&dev_b,
72 N*sizeof(float) ) );
73 HANDLE_ERROR( cudaMalloc( (void**)&dev_partial_c,
74 blocksPerGrid*sizeof(float) ) );
75
76 // fill in the host memory with data
77 for (int i=0; i<N; i++) {
78 a[i] = i;
79 b[i] = i*2;
80 }
81
82 // copy the arrays ‘a‘ and ‘b‘ to the GPU
83 HANDLE_ERROR( cudaMemcpy( dev_a, a, N*sizeof(float),
84 cudaMemcpyHostToDevice ) );
85 HANDLE_ERROR( cudaMemcpy( dev_b, b, N*sizeof(float),
86 cudaMemcpyHostToDevice ) );
87
88 dot<<<blocksPerGrid,threadsPerBlock>>>( dev_a, dev_b,
89 dev_partial_c );
90
91 // copy the array ‘c‘ back from the GPU to the CPU
92 HANDLE_ERROR( cudaMemcpy( partial_c, dev_partial_c,
93 blocksPerGrid*sizeof(float),
94 cudaMemcpyDeviceToHost ) );
95
96 // finish up on the CPU side
97 c = 0;
98 for (int i=0; i<blocksPerGrid; i++) {
99 c += partial_c[i];
100 }
101
102 #define sum_squares(x) (x*(x+1)*(2*x+1)/6)
103 printf( "Does GPU value %.6g = %.6g?\n", c,
104 2 * sum_squares( (float)(N - 1) ) );
105
106 // free memory on the gpu side
107 HANDLE_ERROR( cudaFree( dev_a ) );
108 HANDLE_ERROR( cudaFree( dev_b ) );
109 HANDLE_ERROR( cudaFree( dev_partial_c ) );
110
111 // free memory on the cpu side
112 free( a );
113 free( b );
114 free( partial_c );
115 }
我们首先关注核函数dot。__shared__ float cache[threadsPerBlock];就是这节重点，申请cache数组时，由于使用了共享内存，则每一个block里面都有一份cache，使得block内的thread都可以访问和操作其各自的cache数组。
1 while (tid < N) {
2 temp += a[tid] * b[tid];
3 tid += blockDim.x * gridDim.x;
4 }
这一段我们相当熟悉，每个线程计算若干对a，b的乘积，然后相加。然后这样cache[cacheIndex] = temp;将结果存入cache中。这时，每一个线程的结果都被存在了cache数组中，我们知道接下来要对数组求和，然而这里有潜在的危险，那就是我们不知道所有线程是否已经将数据写入了cache，也就是说，是否每一个线程都已经执行完了第39行。这里就需要等待，等待所有线程执行到同一位置，这就是 __syncthreads();的作用。这个函数称为同步函数，即在所有线程全部执行到__syncthreads()为止，谁也不许动，其后任何代码都无法执行。
因此，我们可以很清楚的明白所有线程全部执行完了第39行，然后同步解除，大家再一起往前走。做加法。
1 int i = blockDim.x/2;
2 while (i != 0) {
3 if (cacheIndex < i)
4 cache[cacheIndex] += cache[cacheIndex + i];
5 __syncthreads();
6 i /= 2;
7 }
8
9 if (cacheIndex == 0)
10 c[blockIdx.x] = cache[0];
这段就不难理解了，逐对相加，最后cache【0】位置的数就是结果。将其值存入c数组，准备导出。
剩下的main函数部分是如下几步操作（和前面学习的差不多）：
1.为输入输出数组分配内存
2.将a，b数组付初值，然后复制给device中，cudaMemcpy（）
3.调用核函数执行并行计算。
4.device值返回后数组c求和。
很明显，由于我们使用了共享内存存储cache数组，使得在操作cache数组时的速度有了大幅提高（相比于全局内存）。共享内存的意义也就在此。
现在，请观察下面的两组代码：
while (i != 0) {
if (cacheIndex < i)
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
i /= 2;
}
while (i != 0) {
if (cacheIndex < i)
{
cache[cacheIndex] += cache[cacheIndex + i];
__syncthreads();
}
i /= 2;
}
下面的代码中由于if的存在，只有部分线程包含同步操作。代码似乎得到了优化。但是真的如此吗
当然不是的，上面的红字“所有线程全部执行到__syncthreads()为止”，所有很重要，<<<>>>中launch了多少个threadperblock，那么就必须要等待所有的线程，一个都不能少。由于if的存在，上例中部分线程永远都不可能执行到cache[cacheIndex] += cache[cacheIndex + i];这一步，因此就要永远等待下去，因而程序无法执行。
总结：在能用共享内存的时候尽量用，进而提高block内的执行效率，但是在同步问题上一定要慎重。。。