这次回顾ECE408 Lecture 12,这次介绍了卷积神经网络。

课程主页:

搬运视频:

卷积层剖析

输入:

  • $A$个形状为$N_1\times N_1$;

卷积层:

  • $B$个形状为$K_1\times K_2$的卷积核;

输出:

  • $A\times B$个输出,形状为$(N_1-K_1+1)\times (N_2-K_2+1)$;

Pooling(降采样)

  • 降采样层
    • 有时内置偏差和非线性;
  • 常见类型
    • 最大值、平均值、L2 范数、加权平均值;
  • 有助于使表示对输入中的大小缩放和平移保持不变;

卷积

前向传播

序列版本的代码

void convLayer_forward(int B, int M, int C, int H, int W, int K, float* X, float* W, float* Y) {
    int H_out = H – K + 1; // calculate H_out, W_out
    int W_out = W – K + 1;
    
    for (int b = 0; b < B; ++b) // for each image
    	for(int m = 0; m < M; m++) // for each output feature map
    		for(int h = 0; h < H_out; h++) // for each output value (two loops)
    			for(int w = 0; w < W_out; w++) {
    				Y[b, m, h, w] = 0.0f; // initialize sum to 0
    					for(int c = 0; c < C; c++) // sum over all input channels
    						for(int p = 0; p < K; p++) // KxK filter
    							for(int q = 0; q < K; q++)
    								Y[b, m, h, w] += X[b, c, h + p, w + q] * W[m, c, p, q];
    			}
}

例子

卷积层中的并行性

输出特征图可以并行计算

  • 通常数量很少,不足以充分利用GPU;

可以并行计算所有输出特征图像素

  • 所有行都可以并行完成;
  • 每行中的所有像素都可以并行完成;
  • 数量很大,但随着我们进入网络的更深层而减少;
  • 所有输入特征图都可以并行处理,但需要原子操作或树归约(稍后学习);
  • 不同的层可能需要不同的策略;

降采样

通过尺度$N$进行降采样

序列版本的代码

void poolingLayer_forward(int B, int M, int H_out, int W_out, int N, float* Y, float* S)
{
	for (int b = 0; b < B; ++b) // for each image
		for (int m = 0; m < M; ++m) // for each output feature map
			for (int x = 0; x < H_out/N; ++x) // for each output value (two loops)
				for (int y = 0; y < W_out/N; ++y) {
					float acc = 0.0f // initialize sum to 0
					for (int p = 0; p < N; ++p) // loop over NxN block of Y (two loops)
						for (int q = 0; q < N; ++q)
							acc += Y[b, m, N*x + p, N*y + q];
					acc /= N * N; // calculate average over block
					S[b, m, x, y] = sigmoid(acc + bias[m]) // bias, non-linearity
				}
}

降采样层的kernel实现

  • 网格到降采样输出特征图像素的直接映射;
  • 在GPU内核中;
    • 需要操作索引映射;
    • 用于访问前一个卷积层的输出特征图像素;
  • 经常合并到前面的卷积层以节省内存带宽;

基本kernel的设计

  • 每个块计算;
    • 一个特征的输出像素块;
    • 每个维度中的$TILE_{width}$个像素;
  • Grid的X维映射到$M$个输出特征图;
  • Grid的Y维度映射到输出特征图中的tiles(线性化顺序);
  • Grid的Z维度用于批量图像,我们从幻灯片中省略了;

例子

假设

  • $M=4$(4 个输出特征图),因此在X维度上有4个块;
  • $W_{out} = H_{out} = 8$(8x8输出特征);

如果$TILE_{WIDTH} = 4$,我们还需要在Y​维度上有4个块:

  • 对于每个输出特征;
  • 每列中的头两个块计算tile的第一行,底部两个计算底部行;

整体的CUDA方法

考虑一个输出特征图:

  • 宽度为$W_out$;
  • 高度是$H_out$;
  • 假设这些是$TILE_{WIDTH}$的整数倍;

令$W_size$为X维度(5)所需的块数,令$H_size$为Y维度(4)中所需的块数。

代码

Host代码:

#define TILE_WIDTH 16 // We will use 4 for small examples.
W_size = W_out/TILE_WIDTH; // number of horizontal tiles per output map
H_size = H_out/TILE_WIDTH; // number of vertical tiles per output map
Y = H_size * W_size;

dim3 blockDim(TILE_WIDTH, TILE_WIDTH, 1); // output tile for untiled code
dim3 gridDim(M, Y, 1);

ConvLayerForward_Kernel<<< gridDim, blockDim >>>();

Kernel代码:

__global__ void ConvLayerForward_Basic_Kernel
	(int C, int W_size, int K, float* X, float* W, float* Y)
{
    int m = blockIdx.x;
    int h = (blockIdx.y / W_size) * TILE_WIDTH + threadIdx.y;
    int w = (blockIdx.y % W_size) * TILE_WIDTH + threadIdx.x;
    float acc = 0.0f;
    for (int c = 0; c < C; c++) { // sum over all input channels
    	for (int p = 0; p < K; p++) // loop over KxK filter
    		for (int q = 0; q < K; q++)
    			acc += X[c, h + p, w + q] * W[m, c, p, q];
    }
    Y[m, h, w] = acc;
}

一些观察

  • 足够的并行度
    • 如果所有输出特征图的像素总数很大(通常是CNN层的情况);
  • 每个输入tile
    • 加载M次(输出特征的数量),因此在全局内存带宽方面效率不高,但X维度的块调度应该会给缓存带来好处;

使用矩阵乘法实现卷积

每个乘积矩阵元素都是一个输出特征图像素:

例1

例2

效率分析

  • 重复的输入特征在输出映射之间共享
    • 有$H_{out} \times W_{out}$个输出特征图元素;
    • 每个都需要来自输入特征映射的$K\times K$个元素;
    • 因此,对于每个输入特征图,复制后的输入元素总数为$H_{out} \times W_{out} \times K\times K$次;
    • 每个原始输入特征图中的元素总数为$(H_{out}+K-1) \times (W_{out}+K-1)$;

假设:

  • $H_{out}=2, W_{out}=2, K=2, C=3$;
  • 那么矩阵形式的输入元素数量为$3\times 2\times 2\times 2 \times 2=48$;
    • 即之前例子的矩阵元素数量;
  • 和$3\times 3 \times 3$的原始元素数量相比,重复因子为$48/27=1.78$;

矩阵乘法相比原始卷积算法的内存访问效率

  • 假设我们使用tiled二维卷积;
  • 对于输入元素
    • 每个输出图块都有$TILE_{WIDTH}^2$个元素;
    • 每个输入图块都有$(TILE_{WIDTH}+K-1)^2$个元素;
    • 输入元素访问总数为$TILE_{WIDTH}^2\times K^2$;
    • 分块算法的缩减系数为$K^2\times TILE_{WIDTH}^2/(TILE_{WIDTH}+K-1)^2$;
  • 卷积滤波器权重元素在每个输出块中重复使用;

Unrolled矩阵的性质

  • 每个展开的列对应一个输出特征图元素;
    • 对于输出特征元素$(h,w)$,展开列的索引为$h\times W_{out}+w$(输出特征图元素的线性化索引);
  • 展开列的每个部分对应一个输入特征图;
    • 展开列的每个部分都有$k\times k$个元素(卷积掩码大小);
    • 对于输入特征图$c$,其在展开列中的部分的垂直索引为$c\times k\times k$(输出特征图元素的线性化索引);

找到输入元素

  • 对于输出元素$(h,w)$,输入特征图$c$的基本索引是$(c, h, w)$;
  • 与卷积掩码元素$(p, q)$相乘的输入元素索引为$(c, h+p, w+q)$;

输入到Unrolled Matrix的代码

Output element (h, w)
Mask element (p, q)
Input feature map c

// calculate the horizontal matrix index
int w_unroll = h * W_out + w;

// find the beginning of the unrolled
int w_base = c * (K*K);

// calculate the vertical matrix index
int h_unroll = w_base + p * K + q;

X_unroll[b, h_unroll, w_unroll] = X[b, c, h + p, w + q];

备注:注意这里$K=k^2$;

完整代码

void unroll(int B, int C, int H, int W, int K, float* X, float* X_unroll)
{
    int H_out = H – K + 1; // calculate H_out, W_out
    int W_out = W – K + 1;
    for (int b = 0; b < B; ++b) // for each image
    	for (int c = 0; c < C; ++c) { // for each input channel
    		int w_base = c * (K*K); // per-channel offset for smallest X_unroll index
    		for (int p = 0; p < K; ++p) // for each element of KxK filter (two loops)
    			for (int q = 0; q < K; ++q) {
    				for (int h = 0; h < H_out; ++h) // for each thread (each output value, two loops)
    					for (int w = 0; w < W_out; ++w) {
    						int h_unroll = w_base + p * K + q; // data needed by one thread
    						int w_unroll = h * W_out + w; // smallest index--across threads (output values)
    						X_unroll[b, h_unroll, w_unroll] = X[b, c, h + p, w + q]; // copy input pixels
    					}
    }
}

小结

卷积层的实现策略

  • 基线
    • Tiled 2D卷积实现,对卷积Mask使用常量内存;
  • 矩阵乘法基线
    • 输入特征图unrolling,卷积Mask的常量内存作为优化;
    • Tiled矩阵乘法kernel;
  • 内置unrolling的矩阵乘法
    • 仅在加载用于矩阵乘法的图块时执行展开;
    • 展开的矩阵只是概念上的;
    • 将概念展开矩阵的tiled元素加载到共享内存时,使用讲座中的属性从输入特征图加载;
  • 更高级的矩阵乘法
    • 使用联合寄存器共享内存tiling;

项目简介

  • 使用CUDA优化修改后的LeNet-5 CNN中卷积层的前向传播(CNN使用Mini-DNN 实现,一种C++框架);
  • 网络将对Fashion MNIST数据集进行分类;
  • 需要注意的一些网络参数
    • 输入大小:86x86 像素,一批10k图像;
    • 输入通道:1;
    • 卷积核大小:7x7;
    • kernel数量:可变(您的代码应该支持这一点);