ECE408 Lecture 12 Convolutional Neural Networks
这次回顾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数量：可变（您的代码应该支持这一点）；
本博客所有文章除特别声明外，均采用 CC BY-NC-SA 4.0 许可协议。转载请注明来自 Doraemonzzz！
相关推荐
评论
ValineLivere