此文是基于个人学习过程随心记录下的，都是个人理解，可能会出现错误。
在cuda_1中讲了host和device的交互方式，还有几个小点要注意下：
- 设备函数即定义只能执行在device上的函数，设备函数只能被kernel和其他设备函数调用。而这个设备函数必须用__device__修饰。
- 用__global__修饰的函数称为核函数，一般由主机调用，在device中执行，但是__global__不可以跟__host__和__device__同时使用。
- host端的普通C++函数可用__host__修饰。对于主机端的函数，__host__修饰符可以省略。可以用__host__和__device__同时修饰一个函数减少冗余，编译器会针对host和device分别编译该函数。
一维模板（1D Stencil） 该窗口或数据范围被依次用于数据集以生成相应结果，我们称为模板操作。为了将一维数组应用与一维模板，得先解释下啥是一维模板：
该模板具有一定宽度，宽度指的是窗口的整体宽度，或者是用于进行逐点计算的基础数据的宽度。这是我们设想一下宽度为7的模板，同时可以定义一个术语--半径。半径就是模板中心点左侧或右侧存在的数据，上图中的半径为3，所以整体宽度为7.我们关注的是一维模板的计算。
如果我们把这个概念带入到block中的话会如何呢（注意此次博客内容大部分都为block层次）
可以看到这是一组数据，我们可以设想这些绿色像素是一组由单个block块处理的数据。我们将输入的数据分配到块，我们使单个块处理的输出元素（以及相应的输入元素）数量等于块的维度（就是块内的线程数，也是blockDim.x）
我们的输入数据为第一个蓝色的block，但是通过一维模板我们得知输入的数据可能不止一个block，而是附带旁边的block
二、通过thread分享data
在block中线程可以通过“共享内存”分享data，因为是片上的memory所以延迟更低带宽更快。所以这叫做共享内存，需要用
__shared__
来启用，需要注意单个block中的共享内存无法在其他block中使用，也就是按块分配的逻辑资源，所以这就是“本地资源”。而全局内存则相反，延迟大但是所有block通用。
那么我们如何使用共享内存来实现一维模板计算？官方slide如下
输入数据集理论上有很多，它代表了我们的一维输入数据集。我们输入数据集中绿色为主要数据集，但是在进行处理的时候还需要橘色的数据集，最后输出的是blockDim.x个数据元素。所以用专业术语就像slide中表示的：
先从全局内存中读blockDim.x + 2*radius个输入元素到共享内存中，然后计算blockDim.x个输出元素，这些输出元素基于共享内存中的数据进行计算的。
最后把blockDim.x个输出元素写回到全局内存中。
所以这时候就明白了一维模板的用处了，每个block都需要一个半径为radius的边界层，进行计算处理的时候需要旁边的边界层数据一起，最后得出结果。其实我个人理解就是有点像卷积运算 ，如果想要让输出数据和输入数据大小一致，那么需要扩充边界才行，并非仅仅输入数据就够。
代码展示：
__global__ void stencil_1d(int *in, int *out) {
//定义静态共享内存 为了存放我们可能用到的元素
__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
//这个是全局唯一thread索引（为了区别上面那行共享内存）
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
//共享内存索引
int lindex = threadIdx.x + RADIUS;
// Read input elements into shared memory（注释1）
temp[lindex] = in[gindex];
if (threadIdx.x < RADIUS) {
//注释2
temp[lindex - RADIUS] = in[gindex - RADIUS];
//注释3
temp[lindex + BLOCK_SIZE] =in[gindex + BLOCK_SIZE];
}
要注意threadIdx.x是独特的线程索引值，每个block内每个thread都有一个独特的threadIdx.x。
注释1:下图中第一行是共享内存未被初始化，也就是全为空的情况。
第二行是 temp <- in[gindex]这一行，虽然没有for循环，但是cuda编程的本质就是每个线程都执行咯。
注释2: 此行对应下图中的第三行，将左边的相关数据加载到共享内存的左边方便后续计算。
注释3:此行对应下图的第四行，将右边的相关数据加载到共享内存的右边方便后续计算。 至此共享内存中的数据全部加载完成。
接下来要对输入数据集的模板宽度进行计算。：
// Apply the stencil
int result = 0;
for (int offset = -RADIUS ; offset <= RADIUS ; offset++)
result += temp[lindex + offset];
// Store the result
out[gindex] = result;
}
每个线程负责从共享内存中读取模板宽度那么长的数据，然后计算结果。这里RADIUS=3，加上输入的thread自己一共为7个thread。所以结果就是7个thread的相加和。
完成该操作后我们利用全局索引将结果值存回全局内存。
但是上述的过程有错误！CUDA编程模型的线程并非同时执行的，线程可以以任意顺序执行。所以如果调用管理线程执行顺序的函数，那么就应该假设线程以任意顺序执行的。
slide又给我们一个最坏场景：
上图说了什么问题呢 ，thread15在thread0获取数据之前已经开始读数据了。要知道输出结果需要类似卷积操作——和周围radius的thread共同计算得出结果，而加载左边数据和右边数据的线程可能并未运行，也就是没初始化。导致加载的数据都是空的，最后得出结果也是错误的。
如何解决这个问题？
线程同步
void __syncthreads( );
线程同步是针对块内的所有线程执行屏障，而非整个网络，这是block内的事情。同步线程并不强制所有块中的所有线程达到某一个特定点，而是在所有线程到达同步线程语句之前，任何线程都不能越过该语句继续执行。但是这是thread级别的。
当使用__syncthreads()时：
- 条件必须在块内保持一致，也就是说所有线程必须能够参与并执行__syncthreads( )。例如我们上面就是每个线程都需要加载共享内存，随后各个线程继续执行其各自的for循环计算result。所以一旦所有线程都加载了共享内存，那么for循环谁先谁后并不会引起我们之前提到了thread竞争的问题了。
- 共享内存的CUDA代码经常会用到此函数，thread将共同加载共享内存，然后在执行任何函数之前用'__synthread()' 来确定共享内存都加载完毕。
所以正确的代码应该是：
__global__ void stencil_1d(int *in, int *out) {
__shared__ int temp[BLOCK_SIZE + 2 * RADIUS];
int gindex = threadIdx.x + blockIdx.x * blockDim.x;
int lindex = threadIdx.x + radius;
// Read input elements intoStencil
temp[lindex] = in[gindex];
if (threadIdx.x < RADIUS) {
temp[lindex – RADIUS] = in[gindex – RADIUS];
temp[lindex + BLOCK_SIZE] = in[gindex + BLOCK_SIZE];
}
// Synchronize (ensure all the data is available)
__syncthreads();
// Apply the stencil
int result = 0;
for (int offset = -RADIUS ; offset <= RADIUS ; offset++)
result += temp[lindex + offset];
// Store the result
out[gindex] = result;
}
在看cuda手册的时候还提到了线程块簇：
- 由多个线程块组成，每个线程块内部都有自己的共享内存。
- 所有线程块的共享内存形成一个分布式共享内存（Distributed Shared Memory）。
所以自下而上的内存层次分为：
- 每线程寄存器和本地内存：每个线程独有的寄存器和本地内存。
- 每块共享内存：每个线程块独有的共享内存。
- 分布式共享内存：每个线程块簇内所有线程块的共享内存合成的分布式共享内存。
- 全局内存：所有GPU内核之间共享的全局内存。
上述总结
上面说了两个知识点，一个是共享内存，一个是线程同步。
- 共享内存用__shared__关键字来分配。同时需要注意共享内存是在块内，对于每个block而言共享内存基本上限制在48Kb（非常小），全局内存基本上以Gb为单位。所以其实共享内存类似于cache，速度快容量小。
- 线程同步是在线程间的同步，也是在block内部，通常和共享内存一同使用。
当一个线程将数据加载到共享内存，而另一个线程使用这些数据的时候我们称之为通信
全局内存中也可以实现通信，但是更常用共享内存。
