使用共享内存优化的GPU加速卷积操作
本教程探讨了使用CUDA在GPU上高效实现卷积操作，重点关注共享内存优化。卷积神经网络（CNN）是现代深度学习在计算机视觉领域的基石，卷积操作占据了它们的大部分计算工作负载。这使其成为GPU加速的主要目标。
您可以在https://github.com/eunomia-bpf/basic-cuda-tutorial找到代码
目录
卷积神经网络简介
卷积神经网络（CNN）通过捕获图像数据中的空间层次结构和模式，彻底改变了计算机视觉领域。它们被设计为通过反向传播自动学习特征的空间层次结构，使用多个构建块，如：
- 卷积层 - 对输入数据应用可学习的滤波器
- 激活函数 - 引入非线性（通常是ReLU）
- 池化层 - 减少空间维度
- 全连接层 - 基于提取的特征执行分类
卷积层是CNN的核心构建块，这就是为什么优化其性能对高效深度学习应用至关重要。
卷积操作
数学定义
2D卷积操作定义为：
其中：
- b 是批次索引
- c 是输入通道索引
- k 是输出通道索引（核数量）
- x、y 是空间坐标
- kx、ky 是核位置
- s 是步长
- p 是填充
维度和内存布局
对于典型的卷积操作：
- 输入形状：[batch_size, in_channels, height, width]
- 核形状：[out_channels, in_channels, kernel_height, kernel_width]
- 输出形状：[batch_size, out_channels, out_height, out_width]
其中：
out_height = (height + 2*padding - kernel_height) / stride + 1
out_width = (width + 2*padding - kernel_width) / stride + 1
实现方法
直接卷积
卷积的朴素实现直接将数学定义映射到代码中：
__global__ void convolutionDirectKernel(
float *input, float *kernels, float *output,
int batchSize, int inputChannels, int inputSize,
int kernelSize, int kernelCount, int outputSize,
int padding, int stride)
{
// 计算输出位置
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int k = blockIdx.z; // 输出通道（核数量）
int b = threadIdx.z; // 批次索引
// 跳过超出范围的线程
if (x >= outputSize || y >= outputSize || k >= kernelCount || b >= batchSize)
return;
// 计算该输出位置的卷积
float sum = 0.0f;
// 对每个输入通道
for (int c = 0; c < inputChannels; c++) {
// 对每个核位置
for (int ky = 0; ky < kernelSize; ky++) {
for (int kx = 0; kx < kernelSize; kx++) {
// 输入位置
int in_x = x * stride - padding + kx;
int in_y = y * stride - padding + ky;
// 如果输入位置在输入范围外则跳过
if (in_x >= 0 && in_x < inputSize && in_y >= 0 && in_y < inputSize) {
// 输入值
float in_val = input[
b * inputChannels * inputSize * inputSize +
c * inputSize * inputSize +
in_y * inputSize + in_x
];
// 核值
float kernel_val = kernels[
k * inputChannels * kernelSize * kernelSize +
c * kernelSize * kernelSize +
ky * kernelSize + kx
];
// 累积结果
sum += in_val * kernel_val;
}
}
}
}
// 存储输出
output[
b * kernelCount * outputSize * outputSize +
k * outputSize * outputSize +
y * outputSize + x
] = sum;
}
直接卷积的特点： - 简单直接的实现 - 每个线程计算一个输出元素 - 全局内存访问冗余度高 - 算术密度低
共享内存优化
优化的关键洞察是相邻的输出元素重用了许多相同的输入值。通过将输入数据加载到共享内存中一次并在多次计算中重用，我们可以显著减少全局内存访问：
__global__ void convolutionSharedKernel(
float *input, float *kernels, float *output,
int batchSize, int inputChannels, int inputSize,
int kernelSize, int kernelCount, int outputSize,
int padding, int stride)
{
// 用于输入瓦片的共享内存
extern __shared__ float sharedData[];
// 计算瓦片尺寸
int tileSize = blockDim.x;
int tileSizeWithPadding = tileSize + kernelSize - 1;
// 线程和块索引
int tx = threadIdx.x;
int ty = threadIdx.y;
int bx = blockIdx.x;
int by = blockIdx.y;
int k = blockIdx.z;
int b = threadIdx.z;
// 将输入数据加载到共享内存
// ...
// 使用共享内存计算卷积
// ...
}
共享内存方法的优势： 1. 减少全局内存访问：每个输入元素只从全局内存加载一次，然后从共享内存中多次重用。 2. 改进内存访问模式：块中的线程访问连续的内存位置。 3. 增加算术密度：每次全局内存访问执行更多计算。
额外的CNN组件
激活函数
激活函数为网络引入非线性。ReLU（修正线性单元）是CNN中最常用的激活函数：
__global__ void reluActivationKernel(float *data, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
data[idx] = fmaxf(0.0f, data[idx]);
}
}
ReLU非常适合并行处理，因为每个输出元素只依赖于单个输入元素。
池化层
池化减少特征图的空间维度，降低计算负载并提供一些平移不变性：
__global__ void maxPoolingKernel(
float *input, float *output,
int batchSize, int channels, int inputSize,
int poolSize, int outputSize, int stride)
{
// ...
// 在池化窗口中找到最大值
for (int dy = 0; dy < poolSize; dy++) {
for (int dx = 0; dx < poolSize; dx++) {
// ...
maxVal = fmaxf(maxVal, value);
}
}
// 存储输出
// ...
}
性能分析
我们的实现比较了两种卷积方法：
- 直接卷积：基准实现，每个线程计算一个输出元素。
- 共享内存卷积：优化实现，将输入瓦片加载到共享内存中。
共享内存优化的典型性能提升： - 对于5×5内核：2-4倍加速 - 对于更大的内核：3-7倍加速 - 对于多输入通道：更大的加速
内存访问分析
对于具有内核大小K×K的直接卷积： - 每个输出元素需要K×K个输入元素 - 对于N×N的输出，需要N×N×K×K次全局内存访问
使用共享内存优化： - 每个输入元素被加载到共享内存一次 - 对于M×M的瓦片（每个维度M个线程），我们加载(M+K-1)×(M+K-1)个元素 - 每个瓦片的全局内存访问总数：(M+K-1)×(M+K-1)
全局内存访问的减少可能相当可观，特别是对于较大的内核大小。
进一步的优化技术
除了本例中展示的共享内存优化外，还有几种其他技术可以进一步加速CNN操作：
-
内核融合：将卷积、偏置加法和激活组合到单个内核中，以减少内核启动开销和内存事务。
-
Winograd算法：减少小内核大小（例如3×3）所需的乘法次数，但代价是增加加法次数。
-
基于FFT的卷积：对于大内核大小，使用快速傅里叶变换可以加速卷积。
-
Im2Col + GEMM：将卷积操作重新格式化为矩阵乘法，以利用高度优化的GEMM库。
-
量化：使用较低精度（INT8、FP16）以增加算术吞吐量并减少内存带宽需求。
-
Tensor Cores：在现代NVIDIA GPU上，利用Tensor Cores进行混合精度矩阵乘法。
-
内核分解：在可能的情况下将较大的内核分解为可分离的1D滤波器（例如，5×5 → 5×1后接1×5）。
结论
卷积操作的高效实现对CNN性能至关重要。通过利用GPU共享内存，我们可以显著减少全局内存访问并提高吞吐量。本例中演示的优化技术代表了现代深度学习框架所基于的基础方法。
对于生产应用，通常建议使用像cuDNN这样的优化库，它们实现了许多这些优化（以及更多），并进行了特定于架构的调整。然而，理解高效卷积的基本原理对于自定义实现和未来的优化是有价值的。