解析：
卷积的物理意义就是加权叠加，常见的操作有一维卷积和二维卷积。需要说明的是，"幽灵元素"通常补0。因为全局存储器和常数存储器变量都是存储在DRAM中的，从DRAM中访问一个变量需要数百甚至数千个时钟周期，从DRAM访问变量的速度通常要比处理器完成算术操作的 速度慢很多。因此，DRAM的长延迟和有限带宽已经成为几乎所有现代处理器的性能瓶颈，通常被称为存储墙问题。为了缓解这个问题，现代处理器通常引进高速 缓存存储器（或多级高速缓存），以减少访问DRAM的次数。
（1）一维卷积
说明：该算法的缺点是受限于访问输入数组d_N的DRAM带宽，掩码数组d_M已经放在L1高速缓存。
（2）分块一维卷积1
我们规定，一个线程块处理的输出元素的集合成为"输出块"；计算不涉及幽灵元素的分块称为"中间块"；输出数据最左边的分块称为左边界分块；输出数据最右边的分块称为右边界分块；涉及多个分块、被多个线程块加载的数据称为"光环元素"或者"边缘元素"。
说明：该算法的优点是将输入数组d_N分块后放入共享存储器，减少了输入数组d_N的DRAM带宽，掩码数组d_M已经放在L1高速缓存。（分块一维卷积1对一维卷积的改进）
（3）分块一维卷积2
说明：该算法的优点是充分利用了L2高速缓存，掩码数组d_M已经放在L1高速缓存。需要说明的是，最近的GPU（比如，Fermi）提供了通用L1和L2高速缓存，L1缓存是每个SM私有的，而L2缓存是所有SM共享的。这样的话，线程块中的光环元素可能存放在L2高速缓存中，我们直接利用光环元素即可，而不再需要将光环元素加载到线程块的共享存储器中。（分块一维卷积2对分块一维卷积1的改进）
2. cudaMemcpyToSymbol
解析：在device中，cudaMemcpyToSymbol进行赋值和读取（__device__，__shared__，__constant__），而在host中，cudaMemcpy进行赋值和读取。需要说明的是，在函数体外声明的变量默认为__device__类型，即全局变量类型。cudaMemcpyToSymbol(d_M, h_M, MaskLen*sizeof(float));和cudaMemcpyFromSymbol(h_M, d_M,MaskLen*sizeof(float));功能相反。常量存储器中的内容通常放在L1高速缓存中，因为常量存储器变量在kernel函数执行期间不会改变，刚好与GPU不提供缓存一致性相符合（主要是为了最大化利用硬件资源，提高处理器的算术运算吞吐率等）。需要说明的是，现代CPU通常都支持处理器核心之间的缓存一致性。
3. shared memory与bank conflict [5][6][7][8] 解析： （1）什么是bank conflict？
在实际中，shared memory被分割成32个等大小的存储体（比如，Maxwell架构），即bank（每个bank拥有每周期32bit=4byte=4char=1int=1float的宽度）。因为一个warp中有32个线程，相当于一个线程对应一个bank。（不同的设备，存储体的数目也不相同，比如设备Tesla架构为16个存储体，后面我们以32个存储体进行讲解）。
特别说明：对于计算能力1.0设备，前个half-warp和后个half-warp不存在bank conflict；对于计算能力2.0设备，前个half-warp和后个half-warp可能存在bank conflict，因为shared memory可以同时让32个bank响应请求。
（2）bank conflict发生的原因？
理想情况下就是不同的线程访问不同的bank，可能是规则的访问，比如，线程0读写bank0，线程1读写bank1，也可能是不规则的，比如线程0读写bank1，线程1读写bank0。这种同一时刻每个bank只被最多1个线程访问的情况下不会出现bank conflict。特殊情况如果有多个线程同时访问同一个bank的同一个地址的时候也不会产生bank conflict，即broadcast。但当多个线程同时访问同一个bank不同的地址时，那么bank conflict就发生了。比如，线程0访问地址0，而线程1访问地址32（归一化后为0），由于它们在同一个bank中，所以就导致了冲突。因为bank conflict发生后，同一个bank的内存读写将被串行化，会导致程序性能大大降低。
（3）什么时候会发生bank conflict呢？
bank conflict主要出现在global memory与shared memory数据交换，及设备函数对shared memory操作中。
（4）如何避免bank conflict呢？
很多时候shared memory的bank conflict可以通过修改数据存储的方式来解决。
4. Thrust [10][11]
Thrust是一款基于GPU CUDA的C++库，其中包含诸多并行算法和数据结构。Thrust主要通过管理系统底层的功能，比如memory access（内存获取）和memory allocation（内存分配）来实现加速，使得工程师们在GPU编程的环境下能更focus在算法的设计上。
（1）容器
host_vector：为主机提供的向量类，并且驻留在主机内存中；device_vector：为设备提供的向量类，并且驻留在设备全局内存中。
说明：fill()，copy()，sequence()。
（2）算法
转换(transformation)；规约(reduction)；前缀求和(prefix sum)；再排序(reordering)；排序(sorting)。
（3）迭代器
constant_iterator；counting_iterator；transform_iterator；permutation_iterator；zip_iterator。
5. 事件，流，纹理内存
解析：
（1）事件：CUDA中的事件本质上是一个GPU时间戳，两个事件的时间差就是算法执行时间。
（2）流：CUDA流表示一个GPU操作队列，可以将一个流看做是GPU上的一个任务，不同任务可以并行执行。
（3）纹理内存：专门为那些在内存访问模式中存在大量空间局部性的图形应用程序而设计的只读内存。
参考文献：
[1] 华科并行计算上机作业：http://wenku.baidu.com/link?url=1tWvUvW0t7BnFChxetS_Mr5_pCF_LZHQGLWxN-
ArVVPccOM_VmoTx9IUD76l_rVMP-iPKWI97vn7wa5ZChz59rr4wlur3rL6k3MGB15qF4W
[2] GPUWattch Energy Model Manual：http://www.gpgpu-sim.org/gpuwattch/
[3] 图形学领域的关键算法及源码链接：http://blog.csdn.net/u013476464/article/details/40857873
[4] 拷贝global memory，cudaMemcpyToSymbol和cudaMemcpy函数是否有区
别：http://blog.csdn.net/litdaguang/article/details/45047015
[5] CUDA GPU编程如何避免bank conflict：http://www.th7.cn/Program/c/201512/719448.shtml
[6] CUDA共享内存bank conflict：http://blog.csdn.net/endlch/article/details/47043069
[7] CUDA bank conflict in shared memory：http://bbs.csdn.net/topics/390836540
[8] CUDA Programming Guide之shared memory的Bank
Confict：http://blog.csdn.net/o_oxo_o/article/details/4296281
[9]Parallel_programming_week3.md：https://github.com/mebusy/notes/blob/c278e037aa8a59aa139fc722d01ed41c
f978921d/dev_notes/Parallel_programming_week3.md
[10] Thrust：http://docs.nvidia.com/cuda/thrust/index.html#axzz4H6gsFZs3
[11] Thrust File List：http://thrust.github.io/doc/files.html
