正如我们在第7.1节中提到的,所有输出(P)元素的计算可以在卷积中并行完成。这使得卷积成为并行计算的理想问题。根据我们在矩阵-矩阵乘法方面的经验,我们可以快速编写一个简单的并行卷积内核。为了简单起见,我们将从1D卷积开始。
第一步是定义内核的主要输入参数。我们假设1D卷积内核收到五个参数:指向输入数组N的指针,指向输入掩码M的指针,指向输出数组P的指针,掩码Mask_Width的大小,以及输入和输出数组宽度的大小。因此,我们有以下设置:
__global__ void convolution_1D_basic_kernel(float *N, float *M,
float *P,
int Mask_Width, int Width) {
// kernel body
}
第二步是确定并实现线程到输出元素的映射。由于输出数组是ID,一个简单而好的方法是将线程组织成一个1D网格,并将每个线程放在网格中来计算一个输出元素。读者应该认识到,就输出元素而言,这与向量加法示例的安排相同。因此,我们可以使用以下语句从每个线程的块索引、块尺寸和线程索引中计算输出元素索引:
int i = blockIdx.x*blockDim.x + threadIdx.x;
一旦我们确定了输出元素索引,我们就可以使用输出元素索引的偏移量访问输入N元素和掩码M元素。为了简单起见,我们假设Mask_Width是一个奇数,卷积是对称的,即M