NVIDIA CUDA 高度并行处理器编程（六）：并行模式：卷积

介绍常数存储器和高速缓存

后面几个文章介绍一些重要的并行计算模式。这些模式是很多并行算法的基础。先介绍卷积，卷积广泛应用于信号处理、图像处理和计算机视觉。卷积通常涉及每个元素上的大量算术运算，比如清晰的图像和视频，计算量很大，不过还好每个输出元素的计算都是相互独立的，我们可以应用并行处理。另一方面，输入元素之间有一定程度的数据共享，还存在一些颇具挑战的边界条件，这使得卷积成为复杂分块方法和输入数据分段的一个重要案例。

背景

卷积是一种数组操作，每个输出元素都是周围输入元素的加权总和。权重是由一个输入掩码数组也叫卷积核（convolution kernel）确定的。下面卷积计算都默认padding。

在音频数字信号处理中，输入数据都是一维的。下面两图为一维卷积的实例：

M为卷积核，N为输入数据。

图7.1即计算：
$P[2]=N[0]\ast M[0]+N[1]\ast M[1]+N[2]\ast M[2]+N[3]\ast M[3]+N[4]\ast M[4]=57$

图7.3即计算：
$P[1]=0\ast M[0]+N[0]\ast M[1]+N[1]\ast M[2]+N[2]\ast M[3]+N[3]\ast M[4]=38$

计算p[1]时，由于N[1]左边只有一个元素，缺少的那个元素用 0 来代替。这些缺失的元素通常称为“幽灵元素”，在并行计算中还会引进其他类型的幽灵元素。这些幽灵元素对分块算法的复杂度和效率影响很大。

对于图像处理和计算机视觉来说，输入数据通常都是二维的。下面为二维图像卷积实例：

计算原理如下图：

带有边界条件的卷积，超出范围的元素当0计算：

一维并行卷积

1. 由于卷积计算与顺序无关，所以可以使用并行解决，先声明kernel函数：

__global__ void convolution_1D_basic_kernel(float *N